From 5da80a92d061ff4dbf03e2dca3c466e82de395d9 Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Fri, 3 Apr 2020 09:50:23 +0200 Subject: [PATCH] fix HIP Philox also make sure that the EK uses only 32-bit floats and that it calls its own RNG wrapper and not the LB's --- src/core/curand_wrapper.hpp | 4 ++-- src/core/grid_based_algorithms/electrokinetics_cuda.cu | 6 +++--- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/src/core/curand_wrapper.hpp b/src/core/curand_wrapper.hpp index 522a9a9a9d4..beac7795d4d 100644 --- a/src/core/curand_wrapper.hpp +++ b/src/core/curand_wrapper.hpp @@ -42,8 +42,8 @@ class philox4x32_10_stateless : private rocrand_device::philox4x32_10_engine { __forceinline__ __device__ uint4 curand_Philox4x32_10(uint4 counter, uint2 key) { - philox4x32_10_stateless *e = nullptr; - return (*e)(counter, key); + philox4x32_10_stateless e; + return e(counter, key); } #endif diff --git a/src/core/grid_based_algorithms/electrokinetics_cuda.cu b/src/core/grid_based_algorithms/electrokinetics_cuda.cu index 028170d4743..2dabde95276 100644 --- a/src/core/grid_based_algorithms/electrokinetics_cuda.cu +++ b/src/core/grid_based_algorithms/electrokinetics_cuda.cu @@ -326,7 +326,7 @@ __device__ void ek_diffusion_migration_lbforce_linkcentered_stencil( float agrid_inv = 1.0f / ek_parameters_gpu->agrid; float sqrt2agrid_inv = 1.0f / (sqrtf(2.0f) * ek_parameters_gpu->agrid); - float sqrt2_inv = 1.0f / sqrt(2.0f); + float sqrt2_inv = 1.0f / sqrtf(2.0f); float twoT_inv = 1.0f / (2.0f * ek_parameters_gpu->T); float D_inv = 1.0f / ek_parameters_gpu->D[species_index]; float force_conv = @@ -1365,7 +1365,7 @@ __device__ void ek_add_fluctuations_to_flux(unsigned int index, for (int i = 0; i < 9; i++) { if (i % 4 == 0) { - random_floats = random_wrapper_philox(index, i + 40, philox_counter); + random_floats = ek_random_wrapper_philox(index, i + 40, philox_counter); random = (random_floats.w - 0.5f) * 2.0f; } else if (i % 4 == 1) { random = (random_floats.x - 0.5f) * 2.0f; @@ -1389,7 +1389,7 @@ __device__ void ek_add_fluctuations_to_flux(unsigned int index, powf(2.0f * average_density * diffusion * time_step / (agrid * agrid), 0.5f) * - random * ek_parameters_gpu->fluctuation_amplitude / sqrt(2.0f); + random * ek_parameters_gpu->fluctuation_amplitude / sqrtf(2.0f); fluc *= !(lb_node.boundary[index] || lb_node.boundary[neighborindex[i]]); #ifdef EK_DEBUG