diff --git a/src/main_gpumd/add_efield.cu b/src/main_gpumd/add_efield.cu
new file mode 100644
index 000000000..081842061
--- /dev/null
+++ b/src/main_gpumd/add_efield.cu
@@ -0,0 +1,168 @@
+/*
+ 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 electric field 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 Ex,
+ const double Ey,
+ const double Ez,
+ const double* g_charge,
+ 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];
+ 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;
+ }
+}
+
+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];
+ 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]];
+ add_efield<<<(group_size - 1) / 64 + 1, 64>>>(
+ group_size,
+ group_size_sum,
+ groups[grouping_method_[call]].contents.data(),
+ 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
+ );
+ CUDA_CHECK_KERNEL
+ }
+}
+
+void Add_Efield::parse(const char** param, int num_param, const std::vector& group)
+{
+ printf("Add electric field.\n");
+
+ // check the number of parameters
+ if (num_param != 6 && num_param != 4) {
+ PRINT_INPUT_ERROR("add_efield 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;
+ efield_table_[num_calls_].resize(table_length_[num_calls_] * 3);
+ if (!is_valid_real(param[3], &efield_table_[num_calls_][0])) {
+ PRINT_INPUT_ERROR("Ex should be a number.\n");
+ }
+ if (!is_valid_real(param[4], &efield_table_[num_calls_][1])) {
+ PRINT_INPUT_ERROR("Ey should be a number.\n");
+ }
+ if (!is_valid_real(param[5], &efield_table_[num_calls_][2])) {
+ PRINT_INPUT_ERROR("Ez should be a number.\n");
+ }
+ printf(" Ex = %g V/A.\n", efield_table_[num_calls_][0]);
+ printf(" Ey = %g V/A.\n", efield_table_[num_calls_][1]);
+ printf(" Ez = %g V/A.\n", efield_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_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_efield file should >= 2.\n");
+ } else {
+ printf(" number of values in the add_efield file = %d.\n", table_length_[num_calls_]);
+ }
+
+ 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 electric field components at each step should be 3.");
+ }
+ for (int t = 0; t < 3; ++t) {
+ efield_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_efield cannot be used more than 10 times in one run.");
+ }
+}
+
+void Add_Efield::finalize()
+{
+ num_calls_ = 0;
+}
diff --git a/src/main_gpumd/add_efield.cuh b/src/main_gpumd/add_efield.cuh
new file mode 100644
index 000000000..52807daa7
--- /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_Efield
+{
+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 efield_table_[10];
+ int grouping_method_[10];
+ int group_id_[10];
+};
diff --git a/src/main_gpumd/add_force.cu b/src/main_gpumd/add_force.cu
index 03f942fe5..a47db7f68 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,
@@ -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/main_gpumd/run.cu b/src/main_gpumd/run.cu
index 3f24014e4..fe28eab4f 100644
--- a/src/main_gpumd/run.cu
+++ b/src/main_gpumd/run.cu
@@ -17,6 +17,7 @@
Run simulation according to the inputs in the run.in file.
------------------------------------------------------------------------------*/
+#include "add_efield.cuh"
#include "add_force.cuh"
#include "cohesive.cuh"
#include "electron_stop.cuh"
@@ -259,6 +260,7 @@ void Run::perform_a_run()
electron_stop.compute(time_step, atom);
add_force.compute(step, group, atom);
+ add_efield.compute(step, group, atom);
integrate.compute2(time_step, double(step) / number_of_steps, group, box, atom, thermo);
@@ -314,6 +316,7 @@ void Run::perform_a_run()
electron_stop.finalize();
add_force.finalize();
+ add_efield.finalize();
integrate.finalize();
mc.finalize();
velocity.finalize();
@@ -466,6 +469,8 @@ void Run::parse_one_keyword(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;
};
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