| 46 | #ifdef AMREX_USE_GPU |
| 47 | namespace { |
| 48 | void ResizeRandomSeed (amrex::ULong gpu_seed) |
| 49 | { |
| 50 | BL_PROFILE("ResizeRandomSeed"); |
| 51 | |
| 52 | using namespace amrex; |
| 53 | |
| 54 | DeallocateRandomSeedDevArray(); |
| 55 | |
| 56 | const int N = Gpu::Device::maxBlocksPerLaunch() * AMREX_GPU_MAX_THREADS; |
| 57 | |
| 58 | #ifdef AMREX_USE_SYCL |
| 59 | |
| 60 | rand_engine_descr = new sycl_rng_descr |
| 61 | (Gpu::Device::streamQueue(), sycl::range<1>(N), gpu_seed, 1); |
| 62 | |
| 63 | gpu_rand_generator = new std::remove_pointer_t<decltype(gpu_rand_generator)> |
| 64 | (Gpu::Device::streamQueue(), gpu_seed+1234ULL); |
| 65 | |
| 66 | #elif defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) |
| 67 | |
| 68 | gpu_rand_state = static_cast<randState_t*>(The_Arena()->alloc(N*sizeof(randState_t))); |
| 69 | randState_t* gpu_rand_state_local = gpu_rand_state; |
| 70 | amrex::ParallelFor(N, [=] AMREX_GPU_DEVICE (int idx) noexcept |
| 71 | { |
| 72 | ULong seqstart = static_cast<ULong>(idx) + 10 * static_cast<ULong>(idx); |
| 73 | AMREX_HIP_OR_CUDA( hiprand_init(gpu_seed, seqstart, 0, &gpu_rand_state_local[idx]);, |
| 74 | curand_init(gpu_seed, seqstart, 0, &gpu_rand_state_local[idx]); ) |
| 75 | }); |
| 76 | |
| 77 | #if defined(AMREX_USE_CUDA) |
| 78 | AMREX_CURAND_SAFE_CALL(curandCreateGenerator |
| 79 | (&gpu_rand_generator, CURAND_RNG_PSEUDO_DEFAULT)); |
| 80 | AMREX_CURAND_SAFE_CALL(curandSetPseudoRandomGeneratorSeed |
| 81 | (gpu_rand_generator, gpu_seed+1234ULL)); |
| 82 | #else |
| 83 | AMREX_HIPRAND_SAFE_CALL(hiprandCreateGenerator |
| 84 | (&gpu_rand_generator, HIPRAND_RNG_PSEUDO_DEFAULT)); |
| 85 | AMREX_HIPRAND_SAFE_CALL(hiprandSetPseudoRandomGeneratorSeed |
| 86 | (gpu_rand_generator, gpu_seed+1234ULL)); |
| 87 | #endif |
| 88 | |
| 89 | #endif |
| 90 | |
| 91 | Gpu::synchronize(); |
| 92 | } |
| 93 | } |
| 94 | #endif |
| 95 |
no test coverage detected