reduction of tuples does not compile

[GoogleCodeExporter]

Please post a short code sample which reproduces the problem:

// transform a sequence into a sequence of tuples
struct tupleMaker: public thrust::unary_function<int, thrust::tuple<int, int> >{

    __host__ __device__
        thrust::tuple<int,int> operator() (int idx){
            int temp = idx % 2;
            return thrust::tuple<int,int>(idx, (temp == 0)? -idx : idx);
        }

};

//add tuples together
struct tuplePlus: public thrust::binary_function<thrust::tuple<int, 
int>,thrust::tuple<int,int>,thrust::tuple<int,int> >{

    __host__ __device__
        thrust::tuple<int,int> operator()(thrust::tuple<int,int> x,thrust::tuple<int,int> y)
        {
            return thrust::tuple<int, int>(thrust::get<0>(x) + thrust::get<0>(y), thrust::get<1>(x) + thrust::get<1>(y));
        }
};

//we want to transform a sequence into a sequence of tuples with 
// alternating sign
thrust::counting_iterator<int> count(0);
tuplePlus binop;
typedef thrust::transform_iterator<tupleMaker, thrust::counting_iterator<int> > 
ti;

ti testb(count, tupleMaker() );
ti teste(count + 32, tupleMaker() );
//test the operator and the iterator
thrust::tuple<int,int> init = testb[0];
thrust::tuple<int,int> a(0,1);
thrust::tuple<int,int> b(2,3);
binop(binop(a,b), init);
thrust::tuple<int,int> temp  = binop(testb[0], testb[1]);
    std::cout << thrust::get<0>(testb[4])<< thrust::get<1>(testb[4]) <<std::endl;
    std::cout << thrust::get<0>(temp)<< thrust::get<1>(temp)<<std::endl;

//this causes a compiler error
thrust::reduce(testb, teste, init, binop);

What is the expected output? What do you see instead?

In file included from /tmp/tmpxft_00003191_00000000-1_fwttest.cudafe1.stub.c:2:
/home/yjko/dev/thrust/iterator/detail/transform_iterator.inl:60: error: 
'typedef class thrust::tuple<int, int, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type> 
thrust::detail::transform_iterator_base<tupleMaker, 
thrust::counting_iterator<int, thrust::use_default, thrust::use_default, 
thrust::use_default>, thrust::use_default, thrust::use_default>::reference' is 
private
/tmp/tmpxft_00003191_00000000-1_fwttest.cudafe1.stub.c:103: error: within this 
context
/home/yjko/dev/thrust/iterator/detail/transform_iterator.inl:70: error: 
'typedef class thrust::tuple<int, int, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type> 
thrust::detail::transform_iterator_base<tupleMaker, 
thrust::counting_iterator<int, thrust::use_default, thrust::use_default, 
thrust::use_default>, thrust::use_default, thrust::use_default>::cv_value_type' 
is private
/tmp/tmpxft_00003191_00000000-1_fwttest.cudafe1.stub.c:106: error: within this 
context

What version of Thrust are you using? Which version of nvcc?  Which host
compiler?  On what operating system?

thrust 1.2, nvcc 3.0, g++ 4.3.2-1.1, debian 2.6.32.22.1.amd64-smp

Please provide any additional information below.

Original issue reported on code.google.com by youngjun...@gmail.com on 8 Nov 2010 at 1:38

[GoogleCodeExporter]

Thanks for the report.  Have you tested this program with the latest version of 
Thrust (1.3) [1] and nvcc (3.2 rc2) [2] ?

[1] http://code.google.com/p/thrust/downloads/list
[2] http://developer.nvidia.com/object/cuda_3_2_toolkit_rc.html

Original comment by jaredhoberock on 8 Nov 2010 at 6:49

• Changed state: Accepted

[GoogleCodeExporter]

hey jared,
thanks for looking into it.
i tried thrust 1.3. 
unfortunately i have to request an upgrade for nvcc from our helpdesk. so it 
might take a while until they get it done.

is it in principle possible, to do what i intended, i.e. transform an iterator 
into a tuple and then use reduce on the tuple iterator?
that would be really great, cause then i would save half the reads and 
reduce-invocations.
its a sweet library btw. saves me a lot of headaches - thanks

Original comment by youngjun...@gmail.com on 8 Nov 2010 at 8:15

[GoogleCodeExporter]

Yes, this is intended to work in the manner you're using it.  My guess is that 
we are simply looking at a bug in nvcc bug which has been fixed in the 
meantime.  The clue is that the error occurs within a file generated by nvcc, 
not a user-generated file.

Anyway, let us know if you're able to try later with a newer nvcc.

Original comment by jaredhoberock on 8 Nov 2010 at 8:20

[GoogleCodeExporter]

i tried to compile it now with nvcc 3.2.
the error i get is the following:

/usr/include/surface_functions.h: In function 'void surf1Dread(T*, 
surface<void, 1>, int, int, cudaSurfaceBoundaryMode)':
/usr/include/surface_functions.h:100: error: there are no arguments to 
'__surf1Dreadc1' that depend on a template parameter, so a declaration of 
'__surf1Dreadc1' must be available
/usr/include/surface_functions.h:100: error: (if you use '-fpermissive', G++ 
will accept your code, but allowing the use of an undeclared name is deprecated)
/usr/include/surface_functions.h:101: error: there are no arguments to 
'__surf1Dreads1' that depend on a template parameter, so a declaration of 
'__surf1Dreads1' must be available
/usr/include/surface_functions.h:102: error: there are no arguments to 
'__surf1Dreadu1' that depend on a template parameter, so a declaration of 
'__surf1Dreadu1' must be available
/usr/include/surface_functions.h:103: error: there are no arguments to 
'__surf1Dreadu2' that depend on a template parameter, so a declaration of 
'__surf1Dreadu2' must be available
/usr/include/surface_functions.h:104: error: there are no arguments to 
'__surf1Dreadu4' that depend on a template parameter, so a declaration of 
'__surf1Dreadu4' must be available
/usr/include/surface_functions.h: In function 'void surf2Dread(T*, 
surface<void, 2>, int, int, int, cudaSurfaceBoundaryMode)':
/usr/include/surface_functions.h:460: error: there are no arguments to 
'__surf2Dreadc1' that depend on a template parameter, so a declaration of 
'__surf2Dreadc1' must be available
/usr/include/surface_functions.h:461: error: there are no arguments to 
'__surf2Dreads1' that depend on a template parameter, so a declaration of 
'__surf2Dreads1' must be available
/usr/include/surface_functions.h:462: error: there are no arguments to 
'__surf2Dreadu1' that depend on a template parameter, so a declaration of 
'__surf2Dreadu1' must be available
/usr/include/surface_functions.h:463: error: there are no arguments to 
'__surf2Dreadu2' that depend on a template parameter, so a declaration of 
'__surf2Dreadu2' must be available
/usr/include/surface_functions.h:464: error: there are no arguments to 
'__surf2Dreadu4' that depend on a template parameter, so a declaration of 
'__surf2Dreadu4' must be available
In file included from /tmp/tmpxft_00001e75_00000000-1_fwttest.cudafe1.stub.c:2:
/home/yjko/dev/thrust/iterator/detail/transform_iterator.inl: At global scope:
/home/yjko/dev/thrust/iterator/detail/transform_iterator.inl:60: error: 
'typedef class thrust::tuple<int, int, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type> 
thrust::detail::transform_iterator_base<tupleMaker, 
thrust::counting_iterator<int, thrust::use_default, thrust::use_default, 
thrust::use_default>, thrust::use_default, thrust::use_default>::reference' is 
private
/tmp/tmpxft_00001e75_00000000-1_fwttest.cudafe1.stub.c:39: error: within this 
context
/home/yjko/dev/thrust/iterator/detail/transform_iterator.inl:70: error: 
'typedef class thrust::tuple<int, int, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, 
thrust::null_type, thrust::null_type> 
thrust::detail::transform_iterator_base<tupleMaker, 
thrust::counting_iterator<int, thrust::use_default, thrust::use_default, 
thrust::use_default>, thrust::use_default, thrust::use_default>::cv_value_type' 
is private
/tmp/tmpxft_00001e75_00000000-1_fwttest.cudafe1.stub.c:42: error: within this 
context

Original comment by youngjun...@gmail.com on 9 Nov 2010 at 12:55

[GoogleCodeExporter]

i get rid of the problem with surface_functions.h with -fpermissive, but the 
original error remains.

Original comment by youngjun...@gmail.com on 9 Nov 2010 at 1:22

[GoogleCodeExporter]

Thanks for trying with the newer nvcc.  I can indeed reproduce this problem 
with Thrust 1.3.0.

This is almost certainly a bug inside nvcc, but I've noticed that nvcc does not 
produce this bug with the development version of Thrust [1].  You can use that 
as a workaround if you like.  I'll file a bug with the compiler team in the 
meantime.

[1] http://code.google.com/p/thrust/source/checkout

Original comment by jaredhoberock on 9 Nov 2010 at 7:30

• Added labels: Milestone-Release1.4

[GoogleCodeExporter]

Forwarded as nvbug-756997

Original comment by jaredhoberock on 9 Nov 2010 at 9:56

• Changed state: Forwarded

[GoogleCodeExporter]

Hey Jared,
that did the trick! thanks.

i have two quick questions if you dont mind:
- if i have a device_vector with content {1,2,3,4} and would use a permuatation 
iterator to get a sequence 1,1,2,2,3,3,4,4 would that result into multiple 
reads from device memory or is there some sort of caching going on?

- is it possible/advisable to use a unary function for a transform iterator 
that maintains an array of dynamic length? can i tell thrust to allocate it in 
shared mem?

thanks

Original comment by youngjun...@gmail.com on 15 Nov 2010 at 12:42