From 2f9e6d4566ad2a2024b1d1baad7046b7e322125c Mon Sep 17 00:00:00 2001
From: Trung Nguyen <ndactrung@gmail.com>
Date: Mon, 21 May 2018 23:32:25 -0500
Subject: [PATCH] Fixed bugs in lal_device.cpp with node_names dynamically
allocated and dipole/long/gpu
---
lib/gpu/Opencl.makefile | 25 +++++++++++++++++++++++--
lib/gpu/lal_device.cpp | 2 +-
lib/gpu/lal_dipole_long_lj.cpp | 4 ++--
lib/gpu/lal_dipole_long_lj.cu | 4 ++--
4 files changed, 28 insertions(+), 7 deletions(-)
diff --git a/lib/gpu/Opencl.makefile b/lib/gpu/Opencl.makefile
index bb07151447..a5fcde68f5 100644
--- a/lib/gpu/Opencl.makefile
+++ b/lib/gpu/Opencl.makefile
@@ -66,7 +66,9 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \
$(OBJ_DIR)/lal_coul_debye.o $(OBJ_DIR)/lal_coul_debye_ext.o \
$(OBJ_DIR)/lal_zbl.o $(OBJ_DIR)/lal_zbl_ext.o \
$(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \
- $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o
+ $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.o \
+ $(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \
+ $(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o
KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \
$(OBJ_DIR)/neighbor_cpu_cl.h $(OBJ_DIR)/pppm_cl.h \
@@ -95,7 +97,8 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \
$(OBJ_DIR)/tersoff_mod_cl.h $(OBJ_DIR)/coul_cl.h \
$(OBJ_DIR)/coul_debye_cl.h $(OBJ_DIR)/zbl_cl.h \
$(OBJ_DIR)/lj_cubic_cl.h $(OBJ_DIR)/vashishta_cl.h \
- $(OBJ_DIR)/ufm_cl.h
+ $(OBJ_DIR)/ufm_cl.h $(OBJ_DIR)/dipole_long_lj_cl.h \
+ $(OBJ_DIR)/lj_expand_coul_long_cl.h
OCL_EXECS = $(BIN_DIR)/ocl_get_devices
@@ -588,6 +591,24 @@ $(OBJ_DIR)/lal_ufm.o: $(ALL_H) lal_ufm.h lal_ufm.cpp $(OBJ_DIR)/ufm_cl.h $(OBJ_
$(OBJ_DIR)/lal_ufm_ext.o: $(ALL_H) lal_ufm.h lal_ufm_ext.cpp lal_base_atomic.h
$(OCL) -o $@ -c lal_ufm_ext.cpp -I$(OBJ_DIR)
+$(OBJ_DIR)/dipole_long_lj_cl.h: lal_dipole_long_lj.cu $(PRE1_H)
+ $(BSH) ./geryon/file_to_cstr.sh dipole_long_lj $(PRE1_H) lal_dipole_long_lj.cu $(OBJ_DIR)/dipole_long_lj_cl.h;
+
+$(OBJ_DIR)/lal_dipole_long_lj.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj.cpp $(OBJ_DIR)/dipole_long_lj_cl.h $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lal_base_charge.o
+ $(OCL) -o $@ -c lal_dipole_long_lj.cpp -I$(OBJ_DIR)
+
+$(OBJ_DIR)/lal_dipole_long_lj_ext.o: $(ALL_H) lal_dipole_long_lj.h lal_dipole_long_lj_ext.cpp lal_base_dipole.h
+ $(OCL) -o $@ -c lal_dipole_long_lj_ext.cpp -I$(OBJ_DIR)
+
+$(OBJ_DIR)/lj_expand_coul_long_cl.h: lal_lj_expand_coul_long.cu $(PRE1_H)
+ $(BSH) ./geryon/file_to_cstr.sh lj_expand_coul_long $(PRE1_H) lal_lj_expand_coul_long.cu $(OBJ_DIR)/lj_expand_coul_long_cl.h;
+
+$(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long.cpp $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/lal_base_charge.o
+ $(OCL) -o $@ -c lal_lj_expand_coul_long.cpp -I$(OBJ_DIR)
+
+$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h
+ $(OCL) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR)
+
$(BIN_DIR)/ocl_get_devices: ./geryon/ucl_get_devices.cpp
$(OCL) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_OPENCL $(OCL_LINK)
diff --git a/lib/gpu/lal_device.cpp b/lib/gpu/lal_device.cpp
index c58b484e4e..0ea128a5b3 100644
--- a/lib/gpu/lal_device.cpp
+++ b/lib/gpu/lal_device.cpp
@@ -80,7 +80,7 @@ int DeviceT::init_device(MPI_Comm world, MPI_Comm replica, const int first_gpu,
char node_name[MPI_MAX_PROCESSOR_NAME];
char *node_names = new char[MPI_MAX_PROCESSOR_NAME*_world_size];
MPI_Get_processor_name(node_name,&name_length);
- MPI_Allgather(&node_name,MPI_MAX_PROCESSOR_NAME,MPI_CHAR,&node_names,
+ MPI_Allgather(&node_name,MPI_MAX_PROCESSOR_NAME,MPI_CHAR,&node_names[0],
MPI_MAX_PROCESSOR_NAME,MPI_CHAR,_comm_world);
std::string node_string=std::string(node_name);
diff --git a/lib/gpu/lal_dipole_long_lj.cpp b/lib/gpu/lal_dipole_long_lj.cpp
index 9cad926a2d..251e1def92 100644
--- a/lib/gpu/lal_dipole_long_lj.cpp
+++ b/lib/gpu/lal_dipole_long_lj.cpp
@@ -156,7 +156,7 @@ void DipoleLongLJT::loop(const bool _eflag, const bool _vflag) {
&this->ans->force, &this->ans->engv, &eflag, &vflag,
&ainum, &nbor_pitch, &this->atom->q,
&this->atom->quat, &cutsq, &_cut_coulsq,
- &_qqrd2e, &this->_threads_per_atom);
+ &_qqrd2e, &_g_ewald, &this->_threads_per_atom);
} else {
this->k_pair.set_size(GX,BX);
this->k_pair.run(&this->atom->x, &lj1, &lj3,
@@ -165,7 +165,7 @@ void DipoleLongLJT::loop(const bool _eflag, const bool _vflag) {
&this->ans->engv, &eflag, &vflag, &ainum,
&nbor_pitch, &this->atom->q,
&this->atom->quat, &cutsq, &_cut_coulsq,
- &_qqrd2e, &this->_threads_per_atom);
+ &_qqrd2e, &_g_ewald, &this->_threads_per_atom);
}
this->time_pair.stop();
}
diff --git a/lib/gpu/lal_dipole_long_lj.cu b/lib/gpu/lal_dipole_long_lj.cu
index 4c1cf9352d..f888dece9b 100644
--- a/lib/gpu/lal_dipole_long_lj.cu
+++ b/lib/gpu/lal_dipole_long_lj.cu
@@ -354,7 +354,7 @@ __kernel void k_dipole_long_lj(const __global numtyp4 *restrict x_,
if (eflag>0) {
acctyp e = (acctyp)0.0;
- if (rsq < lj1[mtype].w && factor_coul > (numtyp)0.0) {
+ if (rsq < cut_coulsq && factor_coul > (numtyp)0.0) {
e = qqrd2e*(b0*g0 + b1*g1 + b2*g2);
if (factor_coul < (numtyp)1.0) {
e_coul *= factor_coul;
@@ -608,7 +608,7 @@ __kernel void k_dipole_long_lj_fast(const __global numtyp4 *restrict x_,
if (eflag>0) {
acctyp e = (acctyp)0.0;
- if (rsq < lj1[mtype].w && factor_coul > (numtyp)0.0) {
+ if (rsq < cut_coulsq && factor_coul > (numtyp)0.0) {
e = qqrd2e*(b0*g0 + b1*g1 + b2*g2);
if (factor_coul < (numtyp)1.0) {
e_coul *= factor_coul;
--
GitLab