jianmu
/
lammps
forked from lijiext/lammps
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Updated the kernels for born/coul/long/cs and coul/long/cs gpu styles

This commit is contained in:
Trung Nguyen 2018-06-19 16:22:00 -05:00
parent 70a7b37614
commit 3ae8fdccd8
3 changed files with 12 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
lib/gpu/lal_born_coul_long_cs.cu
View File

@ -121,7 +121,7 @@ __kernel void k_born_coul_long_cs(const __global numtyp4 *restrict x_,
if (factor_coul<(numtyp)1.0) { if (factor_coul<(numtyp)1.0) {
numtyp grij = g_ewald * (r+EPS_EWALD); numtyp grij = g_ewald * (r+EPS_EWALD);
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= (r+EPS_EWALD); prefactor /= (r+EPS_EWALD);
@ -132,7 +132,7 @@ __kernel void k_born_coul_long_cs(const __global numtyp4 *restrict x_,
} else { } else {
numtyp grij = g_ewald * r; numtyp grij = g_ewald * r;
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= r; prefactor /= r;
@ -262,7 +262,7 @@ __kernel void k_born_coul_long_cs_fast(const __global numtyp4 *restrict x_,
if (factor_coul<(numtyp)1.0) { if (factor_coul<(numtyp)1.0) {
numtyp grij = g_ewald * (r+EPS_EWALD); numtyp grij = g_ewald * (r+EPS_EWALD);
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= (r+EPS_EWALD); prefactor /= (r+EPS_EWALD);
@ -271,10 +271,9 @@ __kernel void k_born_coul_long_cs_fast(const __global numtyp4 *restrict x_,
// scaling of the overall force shall be consistent // scaling of the overall force shall be consistent
r2inv = ucl_recip(rsq + EPS_EWALD_SQR); r2inv = ucl_recip(rsq + EPS_EWALD_SQR);
} else { } else {
numtyp grij = g_ewald * r; numtyp grij = g_ewald * r;
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= r; prefactor /= r;

8
lib/gpu/lal_coul_long_cs.cu
View File

@ -205,7 +205,7 @@ __kernel void k_coul_long_cs(const __global numtyp4 *restrict x_,
if (factor_coul<(numtyp)1.0) { if (factor_coul<(numtyp)1.0) {
numtyp grij = g_ewald * (r+EPS_EWALD); numtyp grij = g_ewald * (r+EPS_EWALD);
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= (r+EPS_EWALD); prefactor /= (r+EPS_EWALD);
@ -217,7 +217,7 @@ __kernel void k_coul_long_cs(const __global numtyp4 *restrict x_,
} else { } else {
numtyp grij = g_ewald * r; numtyp grij = g_ewald * r;
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= r; prefactor /= r;
@ -320,7 +320,7 @@ __kernel void k_coul_long_cs_fast(const __global numtyp4 *restrict x_,
if (factor_coul<(numtyp)1.0) { if (factor_coul<(numtyp)1.0) {
numtyp grij = g_ewald * (r+EPS_EWALD); numtyp grij = g_ewald * (r+EPS_EWALD);
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= (r+EPS_EWALD); prefactor /= (r+EPS_EWALD);
@ -331,7 +331,7 @@ __kernel void k_coul_long_cs_fast(const __global numtyp4 *restrict x_,
} else { } else {
numtyp grij = g_ewald * r; numtyp grij = g_ewald * r;
numtyp expm2 = ucl_exp(-grij*grij); numtyp expm2 = ucl_exp(-grij*grij);
numtyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij); acctyp t = ucl_recip((numtyp)1.0 + CS_EWALD_P*grij);
numtyp u = (numtyp)1.0 - t; numtyp u = (numtyp)1.0 - t;
_erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2; _erfc = t * ((numtyp)1.0 + u*(B0+u*(B1+u*(B2+u*(B3+u*(B4+u*B5)))))) * expm2;
prefactor /= r; prefactor /= r;

4
src/GPU/Install.sh
View File

@ -35,6 +35,8 @@ action pair_beck_gpu.cpp
action pair_beck_gpu.h action pair_beck_gpu.h
action pair_born_coul_long_gpu.cpp pair_born_coul_long.cpp action pair_born_coul_long_gpu.cpp pair_born_coul_long.cpp
action pair_born_coul_long_gpu.h pair_born_coul_long.cpp action pair_born_coul_long_gpu.h pair_born_coul_long.cpp
action pair_born_coul_long_cs_gpu.cpp pair_born_coul_long_cs.cpp
action pair_born_coul_long_cs_gpu.h pair_born_coul_long_cs.cpp
action pair_born_coul_wolf_gpu.cpp action pair_born_coul_wolf_gpu.cpp
action pair_born_coul_wolf_gpu.h action pair_born_coul_wolf_gpu.h
action pair_born_gpu.cpp action pair_born_gpu.cpp
@ -55,6 +57,8 @@ action pair_coul_dsf_gpu.cpp
action pair_coul_dsf_gpu.h action pair_coul_dsf_gpu.h
action pair_coul_long_gpu.cpp pair_coul_long.cpp action pair_coul_long_gpu.cpp pair_coul_long.cpp
action pair_coul_long_gpu.h pair_coul_long.cpp action pair_coul_long_gpu.h pair_coul_long.cpp
action pair_coul_long_cs_gpu.cpp pair_coul_long_cs.cpp
action pair_coul_long_cs_gpu.h pair_coul_long_cs.cpp
action pair_dpd_gpu.cpp action pair_dpd_gpu.cpp
action pair_dpd_gpu.h action pair_dpd_gpu.h
action pair_dpd_tstat_gpu.cpp action pair_dpd_tstat_gpu.cpp