CODEWITHSUNDEEP
compiler-errorscudathrust
Patrick Payne
Patrick Payne

Reputation: 13

Thrust Support of Cuda Vector Types

I am currently trying to make use of the thrust::upper_bound function. I am running into an issue with the arguments that I am supplying to the function. I would like to make use of the CUDA vector types, in particular double3, but when I am using this type I am getting several thrust library errors.

The code block that I am running is below: 

/********************************************************************************                                                                            
eos_search_gpu()                                                                                                                                           
purpose        --- kernel to find the upper bound index for the                                                                                            
                 interpolation values                                                                                                                    
arguments --                                                                                                                                               

y              --- input   double3 values for which we are searching                                                                                       
my             --- input   int number of values for which we are searching                                                                                 
x              --- input   double3 array of structs containin the data table                                                                               
                         values for x, y, and f corresponding to structs                                                                                 
                         ".x", ".y", and ".z"                                                                                                            
n              --- input   int number of data values in the table                                                                                          
dim_x          --- input   int number of data values in the x-direcion of table                                                                            
j[]            --- input/output    int[]  array of int'sthat contains                                                                                      
                 the index of the (x,y,f) position of the upper bound                                                                                    


library calls --                                                                                                                                           

  __host__ __device__ ForwardIterator  thrust::upper_bound(                                                                                                  
         const thrust::detail::execution_policy_base<DerivedPolicy>& exec,                                                                               
         ForwardIterator                                             first,                                                                              
         ForwardIterator                                             last,                                                                               
         const LessThanComparable &                                  value                                                                               
         )                                                                                                                                               

 exec         --- the execution policy to use for parallelization                                                                                        
 first        --- the beginning of the ordered sequence                                                                                                  
 last         --- the end of the ordered sequence                                                                                                        
 value        --- the value to be searched.                                                                                                              

 Returns:     the furthermost iterator i, such that value < *i is false                                                                                  


 const detail::seq_t thrust::seq                                                                                                                            
 an execution policy which requires analgorithm invocation to execute                                                                                    
 sequentially in the current thread.                                                                                                                     

 ********************************************************************************/

__global__ void eos_search_gpu(const double3* y, const int my,
                           const double3* x, const int n,
                           const int dim_x, int * j){

    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if ( i < my) {
      const double ptr = thrust::upper_bound(thrust::seq, x[0].y , x[n-1].y, y[i].y);                                                                     
      j[i] = (ptr - x[i].y - 1);

    }
}

The error messages that are displayed are as follows: 

/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: a class or namespace qualified name is required
      detected during:
        instantiation of class "thrust::iterator_traits<T> [with T=double]" 
 /opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
        instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]" 
  /opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
        instantiation of "ForwardIterator     thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
 /opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
        instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]" 
Interpolation_cuda.cu(254): here

/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: global-scope qualifier (leading "::") is not allowed
      detected during:
        instantiation of class "thrust::iterator_traits<T> [with T=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
        instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
        instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
 /opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
        instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]" 
Interpolation_cuda.cu(254): here

/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: expected a ";"
      detected during:
        instantiation of class "thrust::iterator_traits<T> [with T=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
        instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
        instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
        instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]" 
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
        instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]" 
Interpolation_cuda.cu(254): here

I am wondering if thrust supports the use of CUDA vector types or if I am just doing something incorrectly.

Upvotes: 1

Views: 394

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 151799

You need to satisfy all the expected input types to the thrust algorithm. You're not doing that, as nearly every quantity you've defined doesn't match what thrust is expecting.

For starters, we'll need actual iterators. In device code, this means pointers. Thrust needs to be able to dereference the iterator/pointer, and then you have to instruct thrust what to do with that quantity. For that, we'll need an appropriately defined functor. You may wish to read the thrust quick start guide to understand functor definition and usage. Finally, the sensible pointer/iterator here refers to a double3 type, so we'll need to craft nearly everything to work with double3. Note that we need to choose the version of upper_bound that allows for definition of our own custom-defined functor, so we can manipulate double3 quantities (what we get when we dereference the iterators/pointers) correctly.

This may help:

#include <thrust/binary_search.h>
#include <thrust/execution_policy.h>


struct my_comp_functor{
template <typename T>
__host__ __device__
  bool operator()(T &t1, T &t2) {
    return (t1.y < t2.y);}
};

__global__ void eos_search_gpu(const double3* y, const int my,
                           const double3* x, const int n,
                           const int dim_x, int * j, my_comp_functor my_comp){

    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if ( i < my) {
      const double3 *ptr = thrust::upper_bound(thrust::seq, x, x+n, y[i], my_comp);
      j[i] = (ptr[0].y - x[i].y - 1);

    }
}

int main(){

  double3 *d_y, *d_x;
  int *d_j;

  cudaMalloc(&d_y, 1024);
  cudaMalloc(&d_x, 1024);
  cudaMalloc(&d_j, 1024);
  struct my_comp_functor my_obj;
  eos_search_gpu<<<1,1>>>(d_y, 0, d_x, 0, 0, d_j, my_obj);
  cudaDeviceSynchronize();
}

(the above code compiles without compile error for me on CUDA 9.2, but it is obviously not designed to be functional/useful)

At the end, it seems weird to me that you are jamming a double quantity into j[i] (an integer) but its your code.

Also, I may have gotten the ordering wrong in that functor, so it's possible you may need to change < to >.

When you call this kernel, note I have added a parameter; you'll need to instantiate a my_comp_functor object in host code and then pass that to the kernel in the appropriate location.

Finally, it appears you are doing a vectorized search, note that thrust has vectorized searches available that might obviate the need for this kernel.

Upvotes: 1

Related Questions