From d2de26fc189cd97aa65e939e65c5f31a32274239 Mon Sep 17 00:00:00 2001 From: brucefan1983 Date: Sun, 30 Jun 2024 16:04:02 +0800 Subject: [PATCH 1/3] starting point of add_efield --- src/main_gpumd/add_efield.cuh | 38 ++++++++ src/main_gpumd/add_field.cu | 165 ++++++++++++++++++++++++++++++++++ src/main_gpumd/add_force.cu | 2 +- 3 files changed, 204 insertions(+), 1 deletion(-) create mode 100644 src/main_gpumd/add_efield.cuh create mode 100644 src/main_gpumd/add_field.cu diff --git a/src/main_gpumd/add_efield.cuh b/src/main_gpumd/add_efield.cuh new file mode 100644 index 000000000..cf2f33b0f --- /dev/null +++ b/src/main_gpumd/add_efield.cuh @@ -0,0 +1,38 @@ +/* + Copyright 2017 Zheyong Fan, Ville Vierimaa, Mikko Ervasti, and Ari Harju + This file is part of GPUMD. + GPUMD is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + GPUMD is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + You should have received a copy of the GNU General Public License + along with GPUMD. If not, see . +*/ + +#pragma once + +#include + +class Atom; +class Group; + +class Add_Efeild +{ +public: + + void parse(const char** param, int num_param, const std::vector& group); + void compute(const int step, const std::vector& groups, Atom& atom); + void finalize(); + +private: + + int num_calls_ = 0; + int table_length_[10]; + std::vector force_table_[10]; + int grouping_method_[10]; + int group_id_[10]; +}; diff --git a/src/main_gpumd/add_field.cu b/src/main_gpumd/add_field.cu new file mode 100644 index 000000000..2f93ecc4c --- /dev/null +++ b/src/main_gpumd/add_field.cu @@ -0,0 +1,165 @@ +/* + Copyright 2017 Zheyong Fan, Ville Vierimaa, Mikko Ervasti, and Ari Harju + This file is part of GPUMD. + GPUMD is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + GPUMD is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + You should have received a copy of the GNU General Public License + along with GPUMD. If not, see . +*/ + +/*----------------------------------------------------------------------------80 +Add force to a group of atoms. +------------------------------------------------------------------------------*/ + +#include "add_efield.cuh" +#include "model/atom.cuh" +#include "model/group.cuh" +#include "utilities/read_file.cuh" +#include +#include + +static void __global__ +add_efield( + const int group_size, + const int group_size_sum, + const int* g_group_contents, + const double added_fx, + const double added_fy, + const double added_fz, + double* g_fx, + double* g_fy, + double* g_fz) +{ + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid < group_size) { + const int atom_id = g_group_contents[group_size_sum + tid]; + g_fx[atom_id] += added_fx; + g_fy[atom_id] += added_fy; + g_fz[atom_id] += added_fz; + } +} + +void Add_Efeild::compute(const int step, const std::vector& groups, Atom& atom) +{ + for (int call = 0; call < num_calls_; ++call) { + const int step_mod_table_length = step % table_length_[call]; + const float added_fx = force_table_[call][0 * table_length_[call] + step_mod_table_length]; + const float added_fy = force_table_[call][1 * table_length_[call] + step_mod_table_length]; + const float added_fz = force_table_[call][2 * table_length_[call] + step_mod_table_length]; + const int num_atoms_total = atom.force_per_atom.size() / 3; + const int group_size = groups[grouping_method_[call]].cpu_size[group_id_[call]]; + const int group_size_sum = groups[grouping_method_[call]].cpu_size_sum[group_id_[call]]; + add_efield<<<(group_size - 1) / 64 + 1, 64>>>( + group_size, + group_size_sum, + groups[grouping_method_[call]].contents.data(), + added_fx, + added_fy, + added_fz, + atom.force_per_atom.data(), + atom.force_per_atom.data() + num_atoms_total, + atom.force_per_atom.data() + num_atoms_total * 2 + ); + CUDA_CHECK_KERNEL + } +} + +void Add_Efeild::parse(const char** param, int num_param, const std::vector& group) +{ + printf("Add force.\n"); + + // check the number of parameters + if (num_param != 6 && num_param != 4) { + PRINT_INPUT_ERROR("add_force should have 5 or 3 parameters.\n"); + } + + // parse grouping method + if (!is_valid_int(param[1], &grouping_method_[num_calls_])) { + PRINT_INPUT_ERROR("grouping method should be an integer.\n"); + } + if (grouping_method_[num_calls_] < 0) { + PRINT_INPUT_ERROR("grouping method should >= 0.\n"); + } + if (grouping_method_[num_calls_] >= group.size()) { + PRINT_INPUT_ERROR("grouping method should < maximum number of grouping methods.\n"); + } + + // parse group id + if (!is_valid_int(param[2], &group_id_[num_calls_])) { + PRINT_INPUT_ERROR("group id should be an integer.\n"); + } + if (group_id_[num_calls_] < 0) { + PRINT_INPUT_ERROR("group id should >= 0.\n"); + } + if (group_id_[num_calls_] >= group[grouping_method_[num_calls_]].number) { + PRINT_INPUT_ERROR("group id should < maximum number of groups in the grouping method.\n"); + } + + printf( + " for atoms in group %d of grouping method %d.\n", + group_id_[num_calls_], + grouping_method_[num_calls_] + ); + + if (num_param == 6) { + table_length_[num_calls_] = 1; + force_table_[num_calls_].resize(table_length_[num_calls_] * 3); + if (!is_valid_real(param[3], &force_table_[num_calls_][0])) { + PRINT_INPUT_ERROR("fx should be a number.\n"); + } + if (!is_valid_real(param[4], &force_table_[num_calls_][1])) { + PRINT_INPUT_ERROR("fy should be a number.\n"); + } + if (!is_valid_real(param[5], &force_table_[num_calls_][2])) { + PRINT_INPUT_ERROR("fz should be a number.\n"); + } + printf(" fx = %g eV/A.\n", force_table_[num_calls_][0]); + printf(" fy = %g eV/A.\n", force_table_[num_calls_][1]); + printf(" fz = %g eV/A.\n", force_table_[num_calls_][2]); + } else { + std::ifstream input(param[3]); + if (!input.is_open()) { + printf("Failed to open %s.\n", param[3]); + exit(1); + } + + std::vector tokens = get_tokens(input); + if (tokens.size() != 1) { + PRINT_INPUT_ERROR("The first line of the add_force file should have 1 value."); + } + table_length_[num_calls_] = get_int_from_token(tokens[0], __FILE__, __LINE__); + if (table_length_[num_calls_] < 2) { + PRINT_INPUT_ERROR("Number of steps in the add_force file should >= 2.\n"); + } else { + printf(" number of values in the add_force file = %d.\n", table_length_[num_calls_]); + } + + force_table_[num_calls_].resize(table_length_[num_calls_] * 3); + for (int n = 0; n < table_length_[num_calls_]; ++n) { + std::vector tokens = get_tokens(input); + if (tokens.size() != 3) { + PRINT_INPUT_ERROR("Number of force components at each step should be 3."); + } + for (int t = 0; t < 3; ++t) { + force_table_[num_calls_][t * table_length_[num_calls_] + n] = get_double_from_token(tokens[t], __FILE__, __LINE__); + } + } + } + + ++num_calls_; + + if (num_calls_ > 10) { + PRINT_INPUT_ERROR("add_force cannot be used more than 10 times in one run."); + } +} + +void Add_Efeild::finalize() +{ + num_calls_ = 0; +} diff --git a/src/main_gpumd/add_force.cu b/src/main_gpumd/add_force.cu index 03f942fe5..256cb8245 100644 --- a/src/main_gpumd/add_force.cu +++ b/src/main_gpumd/add_force.cu @@ -24,7 +24,7 @@ Add force to a group of atoms. #include #include -void __global__ +static void __global__ add_force( const int group_size, const int group_size_sum, From aebd547241e02a6d100d9e73327f7f9f6c4a8ac6 Mon Sep 17 00:00:00 2001 From: brucefan1983 Date: Sun, 30 Jun 2024 19:01:43 +0800 Subject: [PATCH 2/3] finish add_efield --- src/main_gpumd/add_efield.cuh | 2 +- src/main_gpumd/add_field.cu | 67 ++++++++++++++++++----------------- src/main_gpumd/add_force.cu | 6 ++-- src/model/atom.cuh | 2 ++ src/model/read_xyz.cu | 41 +++++++++++++++------ 5 files changed, 71 insertions(+), 47 deletions(-) diff --git a/src/main_gpumd/add_efield.cuh b/src/main_gpumd/add_efield.cuh index cf2f33b0f..38d3f7dc5 100644 --- a/src/main_gpumd/add_efield.cuh +++ b/src/main_gpumd/add_efield.cuh @@ -32,7 +32,7 @@ private: int num_calls_ = 0; int table_length_[10]; - std::vector force_table_[10]; + std::vector efield_table_[10]; int grouping_method_[10]; int group_id_[10]; }; diff --git a/src/main_gpumd/add_field.cu b/src/main_gpumd/add_field.cu index 2f93ecc4c..f05e2943f 100644 --- a/src/main_gpumd/add_field.cu +++ b/src/main_gpumd/add_field.cu @@ -14,7 +14,7 @@ */ /*----------------------------------------------------------------------------80 -Add force to a group of atoms. +Add electric field to a group of atoms. ------------------------------------------------------------------------------*/ #include "add_efield.cuh" @@ -29,9 +29,10 @@ add_efield( const int group_size, const int group_size_sum, const int* g_group_contents, - const double added_fx, - const double added_fy, - const double added_fz, + const double Ex, + const double Ey, + const double Ez, + const double* g_charge, double* g_fx, double* g_fy, double* g_fz) @@ -39,9 +40,10 @@ add_efield( const int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < group_size) { const int atom_id = g_group_contents[group_size_sum + tid]; - g_fx[atom_id] += added_fx; - g_fy[atom_id] += added_fy; - g_fz[atom_id] += added_fz; + const double charge = g_charge[atom_id]; + g_fx[atom_id] += charge * Ex; + g_fy[atom_id] += charge * Ey; + g_fz[atom_id] += charge * Ez; } } @@ -49,9 +51,9 @@ void Add_Efeild::compute(const int step, const std::vector& groups, Atom& { for (int call = 0; call < num_calls_; ++call) { const int step_mod_table_length = step % table_length_[call]; - const float added_fx = force_table_[call][0 * table_length_[call] + step_mod_table_length]; - const float added_fy = force_table_[call][1 * table_length_[call] + step_mod_table_length]; - const float added_fz = force_table_[call][2 * table_length_[call] + step_mod_table_length]; + const double Ex = efield_table_[call][0 * table_length_[call] + step_mod_table_length]; + const double Ey = efield_table_[call][1 * table_length_[call] + step_mod_table_length]; + const double Ez = efield_table_[call][2 * table_length_[call] + step_mod_table_length]; const int num_atoms_total = atom.force_per_atom.size() / 3; const int group_size = groups[grouping_method_[call]].cpu_size[group_id_[call]]; const int group_size_sum = groups[grouping_method_[call]].cpu_size_sum[group_id_[call]]; @@ -59,9 +61,10 @@ void Add_Efeild::compute(const int step, const std::vector& groups, Atom& group_size, group_size_sum, groups[grouping_method_[call]].contents.data(), - added_fx, - added_fy, - added_fz, + Ex, + Ey, + Ez, + atom.charge.data(), atom.force_per_atom.data(), atom.force_per_atom.data() + num_atoms_total, atom.force_per_atom.data() + num_atoms_total * 2 @@ -72,11 +75,11 @@ void Add_Efeild::compute(const int step, const std::vector& groups, Atom& void Add_Efeild::parse(const char** param, int num_param, const std::vector& group) { - printf("Add force.\n"); + printf("Add electric field.\n"); // check the number of parameters if (num_param != 6 && num_param != 4) { - PRINT_INPUT_ERROR("add_force should have 5 or 3 parameters.\n"); + PRINT_INPUT_ERROR("add_efield should have 5 or 3 parameters.\n"); } // parse grouping method @@ -109,19 +112,19 @@ void Add_Efeild::parse(const char** param, int num_param, const std::vector tokens = get_tokens(input); if (tokens.size() != 1) { - PRINT_INPUT_ERROR("The first line of the add_force file should have 1 value."); + PRINT_INPUT_ERROR("The first line of the add_efield file should have 1 value."); } table_length_[num_calls_] = get_int_from_token(tokens[0], __FILE__, __LINE__); if (table_length_[num_calls_] < 2) { - PRINT_INPUT_ERROR("Number of steps in the add_force file should >= 2.\n"); + PRINT_INPUT_ERROR("Number of steps in the add_efield file should >= 2.\n"); } else { - printf(" number of values in the add_force file = %d.\n", table_length_[num_calls_]); + printf(" number of values in the add_efield file = %d.\n", table_length_[num_calls_]); } - force_table_[num_calls_].resize(table_length_[num_calls_] * 3); + efield_table_[num_calls_].resize(table_length_[num_calls_] * 3); for (int n = 0; n < table_length_[num_calls_]; ++n) { std::vector tokens = get_tokens(input); if (tokens.size() != 3) { - PRINT_INPUT_ERROR("Number of force components at each step should be 3."); + PRINT_INPUT_ERROR("Number of electric field components at each step should be 3."); } for (int t = 0; t < 3; ++t) { - force_table_[num_calls_][t * table_length_[num_calls_] + n] = get_double_from_token(tokens[t], __FILE__, __LINE__); + efield_table_[num_calls_][t * table_length_[num_calls_] + n] = get_double_from_token(tokens[t], __FILE__, __LINE__); } } } @@ -155,7 +158,7 @@ void Add_Efeild::parse(const char** param, int num_param, const std::vector 10) { - PRINT_INPUT_ERROR("add_force cannot be used more than 10 times in one run."); + PRINT_INPUT_ERROR("add_efield cannot be used more than 10 times in one run."); } } diff --git a/src/main_gpumd/add_force.cu b/src/main_gpumd/add_force.cu index 256cb8245..a47db7f68 100644 --- a/src/main_gpumd/add_force.cu +++ b/src/main_gpumd/add_force.cu @@ -49,9 +49,9 @@ void Add_Force::compute(const int step, const std::vector& groups, Atom& { for (int call = 0; call < num_calls_; ++call) { const int step_mod_table_length = step % table_length_[call]; - const float added_fx = force_table_[call][0 * table_length_[call] + step_mod_table_length]; - const float added_fy = force_table_[call][1 * table_length_[call] + step_mod_table_length]; - const float added_fz = force_table_[call][2 * table_length_[call] + step_mod_table_length]; + const double added_fx = force_table_[call][0 * table_length_[call] + step_mod_table_length]; + const double added_fy = force_table_[call][1 * table_length_[call] + step_mod_table_length]; + const double added_fz = force_table_[call][2 * table_length_[call] + step_mod_table_length]; const int num_atoms_total = atom.force_per_atom.size() / 3; const int group_size = groups[grouping_method_[call]].cpu_size[group_id_[call]]; const int group_size_sum = groups[grouping_method_[call]].cpu_size_sum[group_id_[call]]; diff --git a/src/model/atom.cuh b/src/model/atom.cuh index de571e6eb..8c53d061f 100644 --- a/src/model/atom.cuh +++ b/src/model/atom.cuh @@ -25,11 +25,13 @@ public: std::vector cpu_type; std::vector cpu_type_size; std::vector cpu_mass; + std::vector cpu_charge; std::vector cpu_position_per_atom; std::vector cpu_velocity_per_atom; std::vector cpu_atom_symbol; GPU_Vector type; // per-atom type (1 component) GPU_Vector mass; // per-atom mass (1 component) + GPU_Vector charge; // per-atom charge (1 component) GPU_Vector position_per_atom; // per-atom position (3 components) GPU_Vector position_temp; // used to calculated unwrapped_position GPU_Vector unwrapped_position; // unwrapped per-atom position (3 components) diff --git a/src/model/read_xyz.cu b/src/model/read_xyz.cu index fdb2817b2..591d65edc 100644 --- a/src/model/read_xyz.cu +++ b/src/model/read_xyz.cu @@ -182,6 +182,7 @@ static void read_xyz_line_2( Box& box, int& has_velocity_in_xyz, bool& has_mass, + bool& has_charge, int& num_columns, int* property_offset, std::vector& group) @@ -295,8 +296,9 @@ static void read_xyz_line_2( } // properties - std::string property_name[5] = {"species", "pos", "mass", "vel", "group"}; - int property_position[5] = {-1, -1, -1, -1, -1}; // species,pos,mass,vel,group + const int num_properties = 6; + std::string property_name[num_properties] = {"species", "pos", "mass", "charge", "vel", "group"}; + int property_position[num_properties] = {-1, -1, -1, -1, -1, -1}; // species,pos,mass,charge,vel,group for (int n = 0; n < tokens.size(); ++n) { const std::string properties_string = "properties="; if (tokens[n].substr(0, properties_string.length()) == properties_string) { @@ -308,14 +310,14 @@ static void read_xyz_line_2( } std::vector sub_tokens = get_tokens(line); for (int k = 0; k < sub_tokens.size() / 3; ++k) { - for (int prop = 0; prop < 5; ++prop) { + for (int prop = 0; prop < num_properties; ++prop) { if (sub_tokens[k * 3] == property_name[prop]) { property_position[prop] = k; } } } - if (property_position[3] < 0) { + if (property_position[4] < 0) { has_velocity_in_xyz = 0; printf("Do not specify initial velocities here.\n"); } else { @@ -323,19 +325,19 @@ static void read_xyz_line_2( printf("Specify initial velocities here.\n"); } - if (property_position[4] < 0) { + if (property_position[5] < 0) { group.resize(0); printf("Have no grouping method.\n"); } else { int num_of_grouping_methods = - get_int_from_token(sub_tokens[property_position[4] * 3 + 2], __FILE__, __LINE__); + get_int_from_token(sub_tokens[property_position[5] * 3 + 2], __FILE__, __LINE__); group.resize(num_of_grouping_methods); printf("Have %d grouping method(s).\n", num_of_grouping_methods); } for (int k = 0; k < sub_tokens.size() / 3; ++k) { const int tmp_length = get_int_from_token(sub_tokens[k * 3 + 2], __FILE__, __LINE__); - for (int prop = 0; prop < 5; ++prop) { + for (int prop = 0; prop < num_properties; ++prop) { if (k < property_position[prop]) { property_offset[prop] += tmp_length; } @@ -356,6 +358,11 @@ static void read_xyz_line_2( } else { has_mass = true; } + if (property_position[3] < 0) { + has_charge = false; + } else { + has_charge = true; + } } void read_xyz_in_line_3( @@ -363,6 +370,7 @@ void read_xyz_in_line_3( const int N, const int has_velocity_in_xyz, const bool has_mass, + const bool has_charge, const int num_columns, const int* property_offset, int& number_of_types, @@ -370,6 +378,7 @@ void read_xyz_in_line_3( std::vector& cpu_atom_symbol, std::vector& cpu_type, std::vector& cpu_mass, + std::vector& cpu_charge, std::vector& cpu_position_per_atom, std::vector& cpu_velocity_per_atom, std::vector& group) @@ -377,6 +386,7 @@ void read_xyz_in_line_3( cpu_atom_symbol.resize(N); cpu_type.resize(N); cpu_mass.resize(N); + cpu_charge.resize(N, 0.0); cpu_position_per_atom.resize(N * 3); cpu_velocity_per_atom.resize(N * 3); number_of_types = atom_symbols.size(); @@ -419,18 +429,22 @@ void read_xyz_in_line_3( cpu_mass[n] = MASS_TABLE.at(cpu_atom_symbol[n]); } + if (has_charge) { + cpu_charge[n] = get_double_from_token(tokens[property_offset[3]], __FILE__, __LINE__); + } + if (has_velocity_in_xyz) { const double A_per_fs_to_natural = TIME_UNIT_CONVERSION; for (int d = 0; d < 3; ++d) { cpu_velocity_per_atom[n + N * d] = - get_double_from_token(tokens[property_offset[3] + d], __FILE__, __LINE__) * + get_double_from_token(tokens[property_offset[4] + d], __FILE__, __LINE__) * A_per_fs_to_natural; } } for (int m = 0; m < group.size(); ++m) { group[m].cpu_label[n] = - get_int_from_token(tokens[property_offset[4] + m], __FILE__, __LINE__); + get_int_from_token(tokens[property_offset[5] + m], __FILE__, __LINE__); if (group[m].cpu_label[n] < 0 || group[m].cpu_label[n] >= N) { PRINT_INPUT_ERROR("Group label should >= 0 and < N."); } @@ -536,16 +550,18 @@ void initialize_position( atom_symbols = get_atom_symbols(filename_potential); read_xyz_line_1(input, atom.number_of_atoms); - int property_offset[5] = {0, 0, 0, 0, 0}; // species,pos,mass,vel,group + int property_offset[6] = {0, 0, 0, 0, 0, 0}; // species,pos,mass,vel,group int num_columns = 0; bool has_mass = true; - read_xyz_line_2(input, box, has_velocity_in_xyz, has_mass, num_columns, property_offset, group); + bool has_charge = true; + read_xyz_line_2(input, box, has_velocity_in_xyz, has_mass, has_charge, num_columns, property_offset, group); read_xyz_in_line_3( input, atom.number_of_atoms, has_velocity_in_xyz, has_mass, + has_charge, num_columns, property_offset, number_of_types, @@ -553,6 +569,7 @@ void initialize_position( atom.cpu_atom_symbol, atom.cpu_type, atom.cpu_mass, + atom.cpu_charge, atom.cpu_position_per_atom, atom.cpu_velocity_per_atom, group); @@ -584,6 +601,8 @@ void allocate_memory_gpu(std::vector& group, Atom& atom, GPU_Vector Date: Sun, 30 Jun 2024 19:57:23 +0800 Subject: [PATCH 3/3] call add_efield --- src/main_gpumd/{add_field.cu => add_efield.cu} | 6 +++--- src/main_gpumd/add_efield.cuh | 2 +- src/main_gpumd/run.cu | 5 +++++ src/main_gpumd/run.cuh | 2 ++ 4 files changed, 11 insertions(+), 4 deletions(-) rename src/main_gpumd/{add_field.cu => add_efield.cu} (97%) diff --git a/src/main_gpumd/add_field.cu b/src/main_gpumd/add_efield.cu similarity index 97% rename from src/main_gpumd/add_field.cu rename to src/main_gpumd/add_efield.cu index f05e2943f..081842061 100644 --- a/src/main_gpumd/add_field.cu +++ b/src/main_gpumd/add_efield.cu @@ -47,7 +47,7 @@ add_efield( } } -void Add_Efeild::compute(const int step, const std::vector& groups, Atom& atom) +void Add_Efield::compute(const int step, const std::vector& groups, Atom& atom) { for (int call = 0; call < num_calls_; ++call) { const int step_mod_table_length = step % table_length_[call]; @@ -73,7 +73,7 @@ void Add_Efeild::compute(const int step, const std::vector& groups, Atom& } } -void Add_Efeild::parse(const char** param, int num_param, const std::vector& group) +void Add_Efield::parse(const char** param, int num_param, const std::vector& group) { printf("Add electric field.\n"); @@ -162,7 +162,7 @@ void Add_Efeild::parse(const char** param, int num_param, const std::vector& tokens) electron_stop.parse(param, num_param, atom.number_of_atoms, number_of_types); } else if (strcmp(param[0], "add_force") == 0) { add_force.parse(param, num_param, group); + } else if (strcmp(param[0], "add_efield") == 0) { + add_efield.parse(param, num_param, group); } else if (strcmp(param[0], "mc") == 0) { mc.parse_mc(param, num_param, group, atom); } else if (strcmp(param[0], "dftd3") == 0) { diff --git a/src/main_gpumd/run.cuh b/src/main_gpumd/run.cuh index 4670f04e1..d92579060 100644 --- a/src/main_gpumd/run.cuh +++ b/src/main_gpumd/run.cuh @@ -19,6 +19,7 @@ class Force; class Integrate; class Measure; +#include "add_efield.cuh" #include "add_force.cuh" #include "electron_stop.cuh" #include "force/force.cuh" @@ -70,4 +71,5 @@ private: Measure measure; Electron_Stop electron_stop; Add_Force add_force; + Add_Efield add_efield; };