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

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

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