Actual source code: curand2.cu
1: #include <petsc/private/randomimpl.h>
2: #include <thrust/tuple.h>
3: #include <thrust/transform.h>
4: #include <thrust/device_ptr.h>
5: #include <thrust/iterator/counting_iterator.h>
7: #if PetscDefined(USE_COMPLEX)
8: struct complexscalelw {
9: PetscReal rl, rw;
10: PetscReal il, iw;
12: complexscalelw(PetscScalar low, PetscScalar width)
13: {
14: rl = PetscRealPart(low);
15: il = PetscImaginaryPart(low);
16: rw = PetscRealPart(width);
17: iw = PetscImaginaryPart(width);
18: }
20: __host__ __device__ PetscReal operator()(thrust::tuple<PetscReal, size_t> x) { return thrust::get<1>(x) % 2 ? thrust::get<0>(x) * iw + il : thrust::get<0>(x) * rw + rl; }
21: };
22: #endif
24: struct realscalelw {
25: PetscReal l, w;
27: realscalelw(PetscReal low, PetscReal width) : l(low), w(width) { }
29: __host__ __device__ PetscReal operator()(PetscReal x) { return x * w + l; }
30: };
32: PETSC_INTERN PetscErrorCode PetscRandomCurandScale_Private(PetscRandom r, size_t n, PetscReal *val, PetscBool isneg)
33: {
34: PetscFunctionBegin;
35: if (!r->iset) PetscFunctionReturn(PETSC_SUCCESS);
36: if (isneg) { /* complex case, need to scale differently */
37: #if PetscDefined(USE_COMPLEX)
38: thrust::device_ptr<PetscReal> pval = thrust::device_pointer_cast(val);
39: auto zibit = thrust::make_zip_iterator(thrust::make_tuple(pval, thrust::counting_iterator<size_t>(0)));
40: thrust::transform(zibit, zibit + n, pval, complexscalelw(r->low, r->width));
41: #else
42: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_PLIB, "Negative array size %" PetscInt_FMT, (PetscInt)n);
43: #endif
44: } else {
45: PetscReal rl = PetscRealPart(r->low);
46: PetscReal rw = PetscRealPart(r->width);
47: thrust::device_ptr<PetscReal> pval = thrust::device_pointer_cast(val);
48: thrust::transform(pval, pval + n, pval, realscalelw(rl, rw));
49: }
50: PetscFunctionReturn(PETSC_SUCCESS);
51: }