nested thrust::fill does not work for varied input values

Question

I have tested a minimum test code to fill array "c" with each element of array "a"

It shows that when nested thrust::fill is called with a constant input, it fills the input array this input value correctly.

However, if the input value is a varied value i.e. each element of an array of values, it may fill the input array with only one (first or last) value

#include <thrust/inner_product.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>
#include <thrust/random.h>

#include <thrust/execution_policy.h>

#include <iostream>
#include <cmath>
#include <boost/concept_check.hpp>

struct bFuntor 
{
    bFuntor(int* av__, int* cv__, const int& N__) : av_(av__), cv_(cv__), N_(N__) {};

    __host__ __device__
    int operator()(const int& idx)
    {
      thrust::device_ptr<int> cv_dpt = thrust::device_pointer_cast(cv_);
      thrust::device_ptr<int> cv_dpt1 = thrust::device_pointer_cast(cv_+N_);

      thrust::detail::normal_iterator<thrust::device_ptr<int>> c0 = thrust::detail::make_normal_iterator<thrust::device_ptr<int>>(cv_dpt);
      thrust::detail::normal_iterator<thrust::device_ptr<int>> c1 = thrust::detail::make_normal_iterator<thrust::device_ptr<int>>(cv_dpt1);

      // ** this thrust::fill with varied values does not work
      thrust::fill(thrust::device,c0,c1,av_[idx]);

      // ** this thrust::fill with constant works
//       thrust::fill(thrust::device,c0,c1,10);

      printf("fill result:\n");
      for (int i=0; i<N_; i++)
        printf("fill value: %d -> return value: %d \n",av_[idx],cv_[i]);
      printf("\n");

      return cv_dpt[idx];
    }

    int* av_;
    int* cv_;
    int N_;
};

int main(void)
{
      int N = 2;
      std::vector<int> av = {0,1};
      std::vector<int> cv = {-1,-2};

      thrust::device_vector<int> av_d(N);
      thrust::device_vector<int> cv_d(N);
      av_d = av; cv_d = cv; 

      // call with nested manner
      thrust::transform(thrust::counting_iterator<int>(0),
            thrust::counting_iterator<int>(N),
            cv_d.begin(),
            bFuntor(thrust::raw_pointer_cast(av_d.data()),
            thrust::raw_pointer_cast(cv_d.data()),
                  N));    

      return 0;
}

output case of varied input value:

fill result:
fill value: 0 -> return value: 1 
fill value: 1 -> return value: 1 
fill value: 0 -> return value: 1 
fill value: 1 -> return value: 1 

output case of constant input value:

fill result:
fill value: 10 -> return value: 10 
fill value: 10 -> return value: 10 
fill value: 10 -> return value: 10 
fill value: 10 -> return value: 10 

is this thrust's problem? or it is not supposed to use like this?

Answer 1

This is an example of a data race:

int operator()(const int& idx)
{
  thrust::device_ptr<int> cv_dpt = thrust::device_pointer_cast(cv_);
  thrust::device_ptr<int> cv_dpt1 = thrust::device_pointer_cast(cv_+N_);

  thrust::detail::normal_iterator<thrust::device_ptr<int>> c0 = thrust::detail::make_normal_iterator<thrust::device_ptr<int>>(cv_dpt);
  thrust::detail::normal_iterator<thrust::device_ptr<int>> c1 = thrust::detail::make_normal_iterator<thrust::device_ptr<int>>(cv_dpt1);


  thrust::fill(thrust::device,c0,c1,av_[idx]);

  //.....
}

Here, each call to the functor will attempt to fill the same iterator range (c0 to c1) with different values. Obviously that will produce problems when multiple functor calls occur in parallel.