Permutation_iterators and atomicadd in THRUST

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::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())),

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> 
 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,

It should be safe. By the time your functor is called, the addresses of variables have been fully transformed through the permutation iterator.