Using thrust::reverse with make_zip_iterator+make_tuple

Question

I'm experiencing odd behavior while using the thrust::reverse function on a zip_iterator constructed with a thrust::make_zip_iterator( thrust::make_tuple( )) type syntax (see the answer from JackOLantern here for a good example of that combination).

I wish to reverse some arbitrarily-indicated section of multiple device vectors as in the example code below. When I do the reversing in one go by tupling and zipping them together, unexpected behavior ensues. The first half of the range is correctly changed to an inversion of the second half of the range, however, the second half of the range is left unchanged.

I've been using other thrust functions in a similar fashion (sort_by_key, uniqe_by_key, adjacent_difference, etc.) without issue. Am I just executing this incorrectly or is there some reason that this will not work on a fundamental level? A thought I had is that perhaps the zip_iterator is not bidirectional as required for reverse. Is this true? I couldn't find documentation indicating as such.

A workaround is just to reverse the vector individually, which works as shown below. However, I suspect this will be less efficient. Note that in my actual use-case I have vectors with sizes of the order of 10,000 and I'm zipping up anywhere from 3-7 vectors for the operations.

#include <iostream>
#include <ostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/sequence.h>
#include <thrust/reverse.h>

int main(){

    // initial host vectors
    const int N=10;
    thrust::host_vector<int> h1(N);
    thrust::host_vector<float> h2(N);

    // fill them
    thrust::sequence( h1.begin(), h1.end(), 0);
    thrust::sequence( h2.begin(), h2.end(), 10., 0.5);

    // print initial contents
    for (size_t i=0; i<N; i++){
        std::cout << h1[i] << " " << h2[i] << std::endl;
    }

    // transfer to device
    thrust::device_vector<int> d1 = h1;
    thrust::device_vector<float> d2 = h2;

    // what chunk to invert
    int iStart = 3; int iEnd = 8;

    // attempt to reverse middle via zip_iterators
    thrust::reverse(
            thrust::make_zip_iterator( thrust::make_tuple( d1.begin()+iStart, d2.begin()+iStart)),
            thrust::make_zip_iterator( thrust::make_tuple( d1.begin()+iEnd, d2.begin()+iEnd))
            );

    // pull back and write out unexpected ordering
    thrust::host_vector<int> temp1 = d1;
    thrust::host_vector<float> temp2 = d2;
    std::cout << "<==========>" << std::endl;
    for (size_t i=0; i<N; i++){
        std::cout << temp1[i] << " " << temp2[i] << std::endl;
    }

    // reset device variables
    d1 = h1;
    d2 = h2;

    // reverse individually
    thrust::reverse( d1.begin()+iStart, d1.begin()+iEnd);
    thrust::reverse( d2.begin()+iStart, d2.begin()+iEnd);

    // pull back and write out the desired ordering
    temp1 = d1;
    temp2 = d2;
    std::cout << "<==========>" << std::endl;
    for (size_t i=0; i<N; i++){
        std::cout << temp1[i] << " " << temp2[i] << std::endl;
    }

    return 0;
}

Output

0 10
1 10.5
2 11
3 11.5
4 12
5 12.5
6 13
7 13.5
8 14
9 14.5
<==========>
0 10
1 10.5
2 11
7 13.5
6 13
5 12.5
6 13
7 13.5
8 14
9 14.5
<==========>
0 10
1 10.5
2 11
7 13.5
6 13
5 12.5
4 12
3 11.5
8 14
9 14.5

Answer 1

The information from Robert Crovella in the comments combined with the initially given workaround in the initial post appears to answer the question - thus, I will combine them here so the question can be marked as "answered." If others wish to post other solutions, I'm more than willing to look at them and move the "official answer" check mark. That being said...

The solution to the question has two parts:

1. If using an older version of CUDA and upgrading is an option: upgrade to the newest CUDA version and the operation should work (tested to work on CUDA 9.2.148 - thanks Robert!)

2. If unable to upgrade to a newer version of CUDA: apply reverse to the vectors individually to achieve the same result as given in the initial post. The code with only the working solution is copied below for completeness.

   #include <iostream>
   #include <ostream>
   #include <thrust/device_vector.h>
   #include <thrust/host_vector.h>
   #include <thrust/tuple.h>
   #include <thrust/iterator/zip_iterator.h>
   #include <thrust/sequence.h>
   #include <thrust/reverse.h>
   
   int main(){
   
       // initial host vectors
       const int N=10;
       thrust::host_vector<int> h1(N);
       thrust::host_vector<float> h2(N);
   
       // fill them
       thrust::sequence( h1.begin(), h1.end(), 0);
       thrust::sequence( h2.begin(), h2.end(), 10., 0.5);
   
       // print initial contents
       for (size_t i=0; i<N; i++){
           std::cout << h1[i] << " " << h2[i] << std::endl;
       }
   
       // transfer to device
       thrust::device_vector<int> d1 = h1;
       thrust::device_vector<float> d2 = h2;
   
       // what chunk to invert
       int iStart = 3; int iEnd = 8;
   
       // reverse individually
       thrust::reverse( d1.begin()+iStart, d1.begin()+iEnd);
       thrust::reverse( d2.begin()+iStart, d2.begin()+iEnd);
   
       // pull back and write out the desired ordering
       temp1 = d1;
       temp2 = d2;
       std::cout << "<==========>" << std::endl;
       for (size_t i=0; i<N; i++){
           std::cout << temp1[i] << " " << temp2[i] << std::endl;
       }
   
       return 0;
   }