Hello everybody,

i have a question regarding the usage of permutation_iterators and atomicadd. I have a source vector s_vec(size x) and a target vector t_vec(size y) and indices of the source vector where in the target vector the values should be added i_vec(size x).

I solve that by the following code snippet

```
thrust::for_each(
thrust::make_zip_iterator(thrust::make_tuple(thrust::make_permutation_iterator(t_vec.begin(),i_vec.begin()), s_vec.begin())),
thrust::make_zip_iterator(thrust::make_tuple(thrust::make_permutation_iterator(t_vec.begin(),i_vec.end()), s_vec.end())),
tupleAtomicAdd());
```

where double tupleAtomicAdd (which i have from the web) is defined by

```
__device__ double atomicAdd(double* address, double val) {
unsigned long long int* address_as_ull = (unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do { assumed = old;
old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
} while (assumed != old);
return __longlong_as_double(old); }
struct tupleAtomicAdd:thrust::unary_function<thrust::tuple<double&, double>,double>
{__device__
double operator()(cont thrust::tuple<double&,double> &t)
{ const double &t0 =thrust::get<0>(t);
double *p0=const_cast<double*>(&t0);
return atomicAdd(p0,thrust::get<1>t); }
```

My question is simply if this is safe?

Is there a race condition introduced by the usage of the permutation iterator ?

Because somehow the atomicadd has to check if the address is in use or not. But is this the address of the actual number or some temporare “permutation” number?

Best regards,

And THRUST is really cool,

Bye

Franz