Skip to content

Commit ebdfe26

Browse files
authored
Merge pull request #134 from alicevision/cuda/fix_thrust
[cuda] fix CCTAG_NO_THRUST_COPY_IF code
2 parents 19e3b97 + f337381 commit ebdfe26

File tree

1 file changed

+11
-8
lines changed

1 file changed

+11
-8
lines changed

src/cctag/cuda/frame_07d_vote_if.cu

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,9 @@
99

1010
#include <thrust/copy.h>
1111
#include <thrust/device_ptr.h>
12+
#ifdef CCTAG_NO_THRUST_COPY_IF
13+
#include <thrust/host_vector.h>
14+
#endif
1215

1316
#include <iostream>
1417
#include <algorithm>
@@ -60,21 +63,21 @@ bool Frame::applyVoteIf( )
6063
NumVotersIsGreaterEqual select_op( _voters.dev );
6164

6265
#ifdef CCTAG_NO_THRUST_COPY_IF
63-
#
64-
# There are reports that the Thrust::copy_if fails when you generated code with CUDA 7.0 and run it only
65-
# a 2nd gen Maxwell card (e.g. GTX 980 and GTX 980 Ti). Also, the GTX 1080 seems to be quite similar to
66-
# the GTX 980 and may be affected as well.
67-
# The following code moves everything to host before copy_if, which circumvents the problem but is much
68-
# slower. Make sure that the condition for activating it is very strict.
69-
#
66+
//
67+
// There are reports that the Thrust::copy_if fails when you generated code with CUDA 7.0 and run it only
68+
// a 2nd gen Maxwell card (e.g. GTX 980 and GTX 980 Ti). Also, the GTX 1080 seems to be quite similar to
69+
// the GTX 980 and may be affected as well.
70+
// The following code moves everything to host before copy_if, which circumvents the problem but is much
71+
// slower. Make sure that the condition for activating it is very strict.
72+
//
7073
thrust::host_vector<int> input_host(sz);
7174
thrust::host_vector<int> output_host(sz);
7275
thrust::host_vector<int>::iterator output_host_end;
7376

7477
thrust::copy( input_begin, input_end, input_host.begin() );
7578
output_host_end = thrust::copy_if( input_host.begin(), input_host.end(), output_host.begin(), select_op );
7679
thrust::copy( output_host.begin(), output_host_end, output_begin );
77-
sz = output_host_end = output_host.begin();
80+
sz = output_host_end - output_host.begin();
7881
output_end = output_begin + sz;
7982
#else
8083
output_end = thrust::copy_if( input_begin, input_end, output_begin, select_op );

0 commit comments

Comments
 (0)