Skip to content

Commit

Permalink
Merge branch 'python' into remove_manual_memory_lb
Browse files Browse the repository at this point in the history
  • Loading branch information
@jngrad
jngrad committed Jan 29, 2021
2 parents 62d97d1 + 71bdc0c commit 619725f
Show file tree
Hide file tree
Showing 12 changed files with 312 additions and 256 deletions.
1 change: 1 addition & 0 deletions maintainer/configs/maxset.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,7 @@
#define LB_BOUNDARIES_GPU
#define ELECTROKINETICS
#define EK_BOUNDARIES
#define EK_DEBUG
#define MMM1D_GPU
#endif

Expand Down
63 changes: 33 additions & 30 deletions src/core/grid_based_algorithms/electrokinetics_cuda.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -156,6 +156,10 @@ EK_parameters ek_parameters = {
nullptr,
// lb_force_density_previous
nullptr,
#ifdef EK_DEBUG
// j_fluc
nullptr,
#endif
// rho
{nullptr},
// species_index
Expand Down Expand Up @@ -2626,7 +2630,7 @@ LOOKUP_TABLE default\n",
lbpar_gpu.agrid, lbpar_gpu.number_of_nodes);

for (int i = 0; i < lbpar_gpu.number_of_nodes; i++) {
fprintf(fp, "%e %e %e ", host_values[i].v[0] * lattice_speed,
fprintf(fp, "%e %e %e\n", host_values[i].v[0] * lattice_speed,
host_values[i].v[1] * lattice_speed,
host_values[i].v[2] * lattice_speed);
}
Expand Down Expand Up @@ -2689,7 +2693,7 @@ LOOKUP_TABLE default\n",
lbpar_gpu.agrid, lbpar_gpu.number_of_nodes);
auto const agrid = lb_lbfluid_get_agrid();
for (int i = 0; i < lbpar_gpu.number_of_nodes; i++) {
fprintf(fp, "%e ", host_values[i].rho / agrid / agrid / agrid);
fprintf(fp, "%e\n", host_values[i].rho / agrid / agrid / agrid);
}

free(host_values);
Expand Down Expand Up @@ -3239,39 +3243,38 @@ int ek_print_vtk_flux_fluc(int species, char *filename) {
return 1;
}

ekfloat *fluxes = (ekfloat *)Utils::malloc(ek_parameters.number_of_nodes *
13 * sizeof(ekfloat));
if (ek_parameters.species_index[species] == -1) {
return 1;
}

if (ek_parameters.species_index[species] != -1) {
int threads_per_block = 64;
int blocks_per_grid_y = 4;
int blocks_per_grid_x = (ek_parameters.number_of_nodes +
threads_per_block * blocks_per_grid_y - 1) /
(threads_per_block * blocks_per_grid_y);
dim3 dim_grid = make_uint3(blocks_per_grid_x, blocks_per_grid_y, 1);
auto *fluxes = (ekfloat *)Utils::malloc(ek_parameters.number_of_nodes * 13 *
sizeof(ekfloat));

KERNELCALL(ek_clear_fluxes, dim_grid, threads_per_block);
KERNELCALL(ek_calculate_quantities, dim_grid, threads_per_block,
ek_parameters.species_index[species], *current_nodes, node_f,
ek_lbparameters_gpu, ek_lb_device_values,
philox_counter.value());
reset_LB_force_densities_GPU(false);
int threads_per_block = 64;
int blocks_per_grid_y = 4;
int blocks_per_grid_x = (static_cast<int>(ek_parameters.number_of_nodes) +
threads_per_block * blocks_per_grid_y - 1) /
(threads_per_block * blocks_per_grid_y);
dim3 dim_grid = make_uint3(blocks_per_grid_x, blocks_per_grid_y, 1);

KERNELCALL(ek_clear_fluxes, dim_grid, threads_per_block);
KERNELCALL(ek_calculate_quantities, dim_grid, threads_per_block,
ek_parameters.species_index[species], *current_nodes, node_f,
ek_lbparameters_gpu, ek_lb_device_values, philox_counter.value());
reset_LB_force_densities_GPU(false);

#ifdef EK_BOUNDARIES
KERNELCALL(ek_apply_boundaries, dim_grid, threads_per_block,
ek_parameters.species_index[species], *current_nodes, node_f);
KERNELCALL(ek_apply_boundaries, dim_grid, threads_per_block,
ek_parameters.species_index[species], *current_nodes, node_f);
#endif

cuda_safe_mem(
cudaMemcpy(fluxes, ek_parameters.j_fluc,
ek_parameters.number_of_nodes * 13 * sizeof(ekfloat),
cudaMemcpyDeviceToHost));
} else
return 1;
cuda_safe_mem(cudaMemcpy(fluxes, ek_parameters.j_fluc,
ek_parameters.number_of_nodes * 13 * sizeof(ekfloat),
cudaMemcpyDeviceToHost));

fprintf(fp, "\
# vtk DataFile Version 2.0\n\
flux_%d\n\
flux_fluc_%d\n\
ASCII\n\
\n\
DATASET STRUCTURED_POINTS\n\
Expand All @@ -3280,7 +3283,7 @@ ORIGIN %f %f %f\n\
SPACING %f %f %f\n\
\n\
POINT_DATA %u\n\
SCALARS flux_%d float 3\n\
SCALARS flux_fluc_%d float 4\n\
LOOKUP_TABLE default\n",
species, ek_parameters.dim_x, ek_parameters.dim_y,
ek_parameters.dim_z, ek_parameters.agrid * 0.5f,
Expand Down Expand Up @@ -3507,7 +3510,7 @@ int ek_print_vtk_flux_link(int species, char *filename) {

fprintf(fp, "\
# vtk DataFile Version 2.0\n\
flux_%d\n\
flux_link_%d\n\
ASCII\n\
\n\
DATASET STRUCTURED_POINTS\n\
Expand All @@ -3516,7 +3519,7 @@ ORIGIN %f %f %f\n\
SPACING %f %f %f\n\
\n\
POINT_DATA %u\n\
SCALARS flux_%d float 3\n\
SCALARS flux_link_%d float 13\n\
LOOKUP_TABLE default\n",
species, ek_parameters.dim_x, ek_parameters.dim_y,
ek_parameters.dim_z, ek_parameters.agrid * 0.5f,
Expand All @@ -3527,7 +3530,7 @@ LOOKUP_TABLE default\n",
for (int i = 0; i < ek_parameters.number_of_nodes; i++) {
rhoindex_linear2cartesian_host(i, coord);

fprintf(fp, "%e %e %e %e %e %e %e %e %e %e %e %e %e \n",
fprintf(fp, "%e %e %e %e %e %e %e %e %e %e %e %e %e\n",
fluxes[jindex_getByRhoLinear_host(i, 0)],
fluxes[jindex_getByRhoLinear_host(i, 1)],
fluxes[jindex_getByRhoLinear_host(i, 2)],
Expand Down
2 changes: 1 addition & 1 deletion src/core/grid_based_algorithms/lb.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -998,6 +998,7 @@ void lattice_boltzmann_update() {
/** \name Coupling part */
/***********************************************************************/
/**@{*/
#ifdef ADDITIONAL_CHECKS
template <class T> int compare_buffers(T const &buff_a, T const &buff_b) {
if (buff_a != buff_b) {
runtimeErrorMsg() << "Halo buffers are not identical";
Expand All @@ -1006,7 +1007,6 @@ template <class T> int compare_buffers(T const &buff_a, T const &buff_b) {
return ES_OK;
}

#ifdef ADDITIONAL_CHECKS
/** Check consistency of the halo regions.
* Test whether the halo regions have been exchanged correctly.
*/
Expand Down
144 changes: 97 additions & 47 deletions src/python/espressomd/electrokinetics.pxd
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,53 +28,103 @@ IF ELECTROKINETICS and CUDA:
DEF MAX_NUMBER_OF_SPECIES = 10

# EK data struct
ctypedef struct EK_parameters:
float agrid
float time_step
float lb_density
unsigned int dim_x
unsigned int dim_y
unsigned int dim_z
unsigned int number_of_nodes
float viscosity
float bulk_viscosity
float gamma_odd
float gamma_even
float friction
float T
float prefactor
float lb_force_density[3]
unsigned int number_of_species
int reaction_species[3]
float rho_reactant_reservoir
float rho_product0_reservoir
float rho_product1_reservoir
float reaction_ct_rate
float reaction_fraction_0
float reaction_fraction_1
float mass_reactant
float mass_product0
float mass_product1
int stencil
int number_of_boundary_nodes
float fluctuation_amplitude
bool fluctuations
bool advection
bool fluidcoupling_ideal_contribution
float * charge_potential
ekfloat * j
float * lb_force_density_previous
ekfloat * rho[MAX_NUMBER_OF_SPECIES]
int species_index[MAX_NUMBER_OF_SPECIES]
float density[MAX_NUMBER_OF_SPECIES]
float D[MAX_NUMBER_OF_SPECIES]
float d[MAX_NUMBER_OF_SPECIES]
float valency[MAX_NUMBER_OF_SPECIES]
float ext_force_density[3][MAX_NUMBER_OF_SPECIES]
char * node_is_catalyst
bool es_coupling
float * charge_potential_buffer
float * electric_field
IF EK_DEBUG:
ctypedef struct EK_parameters:
float agrid
float time_step
float lb_density
unsigned int dim_x
unsigned int dim_y
unsigned int dim_z
unsigned int number_of_nodes
float viscosity
float bulk_viscosity
float gamma_odd
float gamma_even
float friction
float T
float prefactor
float lb_force_density[3]
unsigned int number_of_species
int reaction_species[3]
float rho_reactant_reservoir
float rho_product0_reservoir
float rho_product1_reservoir
float reaction_ct_rate
float reaction_fraction_0
float reaction_fraction_1
float mass_reactant
float mass_product0
float mass_product1
int stencil
int number_of_boundary_nodes
float fluctuation_amplitude
bool fluctuations
bool advection
bool fluidcoupling_ideal_contribution
float * charge_potential
ekfloat * j
float * lb_force_density_previous
ekfloat * j_fluc
ekfloat * rho[MAX_NUMBER_OF_SPECIES]
int species_index[MAX_NUMBER_OF_SPECIES]
float density[MAX_NUMBER_OF_SPECIES]
float D[MAX_NUMBER_OF_SPECIES]
float d[MAX_NUMBER_OF_SPECIES]
float valency[MAX_NUMBER_OF_SPECIES]
float ext_force_density[3][MAX_NUMBER_OF_SPECIES]
char * node_is_catalyst
bool es_coupling
float * charge_potential_buffer
float * electric_field
ELSE:
ctypedef struct EK_parameters:
float agrid
float time_step
float lb_density
unsigned int dim_x
unsigned int dim_y
unsigned int dim_z
unsigned int number_of_nodes
float viscosity
float bulk_viscosity
float gamma_odd
float gamma_even
float friction
float T
float prefactor
float lb_force_density[3]
unsigned int number_of_species
int reaction_species[3]
float rho_reactant_reservoir
float rho_product0_reservoir
float rho_product1_reservoir
float reaction_ct_rate
float reaction_fraction_0
float reaction_fraction_1
float mass_reactant
float mass_product0
float mass_product1
int stencil
int number_of_boundary_nodes
float fluctuation_amplitude
bool fluctuations
bool advection
bool fluidcoupling_ideal_contribution
float * charge_potential
ekfloat * j
float * lb_force_density_previous
ekfloat * rho[MAX_NUMBER_OF_SPECIES]
int species_index[MAX_NUMBER_OF_SPECIES]
float density[MAX_NUMBER_OF_SPECIES]
float D[MAX_NUMBER_OF_SPECIES]
float d[MAX_NUMBER_OF_SPECIES]
float valency[MAX_NUMBER_OF_SPECIES]
float ext_force_density[3][MAX_NUMBER_OF_SPECIES]
char * node_is_catalyst
bool es_coupling
float * charge_potential_buffer
float * electric_field

cdef extern EK_parameters ek_parameters

Expand Down
6 changes: 3 additions & 3 deletions src/python/espressomd/electrokinetics.pyx
View file
Original file line number Diff line number Diff line change
Expand Up @@ -366,7 +366,7 @@ IF ELECTROKINETICS:
self.node[1] = key[1]
self.node[2] = key[2]
if not lb_lbnode_is_index_valid(self.node):
raise ValueError("LB node index out of bounds")
raise IndexError("LB node index out of bounds")

property potential:
def __get__(self):
Expand Down Expand Up @@ -544,7 +544,7 @@ IF ELECTROKINETICS:
self.node[2] = key[2]
self.id = id
if not lb_lbnode_is_index_valid(self.node):
raise ValueError("LB node index out of bounds")
raise IndexError("LB node index out of bounds")

property density:
def __set__(self, value):
Expand Down Expand Up @@ -572,4 +572,4 @@ IF ELECTROKINETICS:
self.id, self.node[0], self.node[1], self.node[2], flux) != 0:
raise Exception("Species has not been added to EK.")

return np.array(flux[0], flux[1], flux[2])
return np.array([flux[0], flux[1], flux[2]])
10 changes: 3 additions & 7 deletions testsuite/python/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -125,9 +125,9 @@ python_test(FILE lb_stokes_sphere.py MAX_NUM_PROC 4 LABELS gpu long)
python_test(FILE lb_pressure_tensor.py MAX_NUM_PROC 1 LABELS gpu long)
python_test(FILE ek_fluctuations.py MAX_NUM_PROC 1 LABELS gpu)
python_test(FILE ek_charged_plate.py MAX_NUM_PROC 1 LABELS gpu)
python_test(FILE ek_eof_one_species_x.py MAX_NUM_PROC 1 LABELS gpu)
python_test(FILE ek_eof_one_species_y.py MAX_NUM_PROC 1 LABELS gpu)
python_test(FILE ek_eof_one_species_z.py MAX_NUM_PROC 1 LABELS gpu)
python_test(FILE ek_eof_one_species.py MAX_NUM_PROC 1 LABELS gpu SUFFIX x)
python_test(FILE ek_eof_one_species.py MAX_NUM_PROC 1 LABELS gpu SUFFIX y)
python_test(FILE ek_eof_one_species.py MAX_NUM_PROC 1 LABELS gpu SUFFIX z)
python_test(FILE exclusions.py MAX_NUM_PROC 2)
python_test(FILE langevin_thermostat.py MAX_NUM_PROC 1)
python_test(FILE langevin_thermostat_stats.py MAX_NUM_PROC 1 LABELS long)
Expand Down Expand Up @@ -248,10 +248,6 @@ add_custom_target(
${CMAKE_CURRENT_BINARY_DIR}
COMMAND
${CMAKE_COMMAND} -E copy ${CMAKE_CURRENT_SOURCE_DIR}/thermostats_common.py
${CMAKE_CURRENT_BINARY_DIR}
COMMAND
${CMAKE_COMMAND} -E copy
${CMAKE_CURRENT_SOURCE_DIR}/ek_eof_one_species_base.py
${CMAKE_CURRENT_BINARY_DIR})

add_custom_target(
Expand Down
Loading

0 comments on commit 619725f

Please sign in to comment.