Compare commits
51 Commits
patch_26Ja
...
patch_13Fe
| Author | SHA1 | Date | |
|---|---|---|---|
| cb982f2f28 | |||
| 4843296d4e | |||
| 2bdda8f6c0 | |||
| 0068ef5616 | |||
| 02b0e6cc55 | |||
| fbb24c2406 | |||
| a5f830c40c | |||
| 8c074a363a | |||
| 27aca14094 | |||
| 191453e1c7 | |||
| 207adc3968 | |||
| 84c517159d | |||
| 6ca377436f | |||
| dc34a32602 | |||
| 067119f6c6 | |||
| 1834a5e46c | |||
| 6a4918b39a | |||
| 5da0d39392 | |||
| 6f92429602 | |||
| 38e0e4bb69 | |||
| daf9f95381 | |||
| 6595fde0a1 | |||
| 6bcec9c61d | |||
| 9d1991bf84 | |||
| 0a87b7443a | |||
| 7ee45ec5f3 | |||
| d4c9e2500b | |||
| 6232073d3b | |||
| ed59193d13 | |||
| 67bed8e853 | |||
| bcb1d94b9a | |||
| fbe30b5683 | |||
| 9ef55fedf7 | |||
| 997142a4c1 | |||
| 033b07fdb7 | |||
| 51a0b6b445 | |||
| 59f4a77dd5 | |||
| 579cc6d7aa | |||
| 5afd3e995b | |||
| 2a6f5e651c | |||
| 09fc8b0bd7 | |||
| e5d0bde783 | |||
| 9daf7fb650 | |||
| b5d622c6a3 | |||
| 2023fa28e0 | |||
| 5b29515849 | |||
| 5b18421dd2 | |||
| cf95ea0709 | |||
| 6a74a81da0 | |||
| f0a4ed615d | |||
| cfe818a175 |
@ -1,7 +1,7 @@
|
||||
<!-- HTML_ONLY -->
|
||||
<HEAD>
|
||||
<TITLE>LAMMPS Users Manual</TITLE>
|
||||
<META NAME="docnumber" CONTENT="26 Jan 2017 version">
|
||||
<META NAME="docnumber" CONTENT="13 Feb 2017 version">
|
||||
<META NAME="author" CONTENT="http://lammps.sandia.gov - Sandia National Laboratories">
|
||||
<META NAME="copyright" CONTENT="Copyright (2003) Sandia Corporation. This software and manual is distributed under the GNU General Public License.">
|
||||
</HEAD>
|
||||
@ -21,7 +21,7 @@
|
||||
<H1></H1>
|
||||
|
||||
LAMMPS Documentation :c,h3
|
||||
26 Jan 2017 version :c,h4
|
||||
13 Feb 2017 version :c,h4
|
||||
|
||||
Version info: :h4
|
||||
|
||||
|
||||
@ -969,7 +969,7 @@ KOKKOS, o = USER-OMP, t = OPT.
|
||||
"lubricateU/poly"_pair_lubricateU.html,
|
||||
"meam"_pair_meam.html,
|
||||
"mie/cut (o)"_pair_mie.html,
|
||||
"morse (got)"_pair_morse.html,
|
||||
"morse (gkot)"_pair_morse.html,
|
||||
"nb3b/harmonic (o)"_pair_nb3b_harmonic.html,
|
||||
"nm/cut (o)"_pair_nm.html,
|
||||
"nm/cut/coul/cut (o)"_pair_nm.html,
|
||||
|
||||
@ -22,7 +22,7 @@ either conceptually, or as printed out by the program.
|
||||
|
||||
12.1 Common problems :link(err_1),h4
|
||||
|
||||
If two LAMMPS runs do not produce the same answer on different
|
||||
If two LAMMPS runs do not produce the exact same answer on different
|
||||
machines or different numbers of processors, this is typically not a
|
||||
bug. In theory you should get identical answers on any number of
|
||||
processors and on any machine. In practice, numerical round-off can
|
||||
@ -80,12 +80,24 @@ order. If you mess this up, LAMMPS will often flag the error, but it
|
||||
may also simply read a bogus argument and assign a value that is
|
||||
valid, but not what you wanted. E.g. trying to read the string "abc"
|
||||
as an integer value of 0. Careful reading of the associated doc page
|
||||
for the command should allow you to fix these problems. Note that
|
||||
some commands allow for variables to be specified in place of numeric
|
||||
constants so that the value can be evaluated and change over the
|
||||
course of a run. This is typically done with the syntax {v_name} for
|
||||
a parameter, where name is the name of the variable. This is only
|
||||
allowed if the command documentation says it is.
|
||||
for the command should allow you to fix these problems. In most cases,
|
||||
where LAMMPS expects to read a number, either integer or floating point,
|
||||
it performs a stringent test on whether the provided input actually
|
||||
is an integer or floating-point number, respectively, and reject the
|
||||
input with an error message (for instance, when an integer is required,
|
||||
but a floating-point number 1.0 is provided):
|
||||
|
||||
ERROR: Expected integer parameter in input script or data file :pre
|
||||
|
||||
Some commands allow for using variable references in place of numeric
|
||||
constants so that the value can be evaluated and may change over the
|
||||
course of a run. This is typically done with the syntax {v_name} for a
|
||||
parameter, where name is the name of the variable. On the other hand,
|
||||
immediate variable expansion with the syntax ${name} is performed while
|
||||
reading the input and before parsing commands,
|
||||
|
||||
NOTE: Using a variable reference (i.e. {v_name}) is only allowed if
|
||||
the documentation of the corresponding command explicitly says it is.
|
||||
|
||||
Generally, LAMMPS will print a message to the screen and logfile and
|
||||
exit gracefully when it encounters a fatal error. Sometimes it will
|
||||
|
||||
@ -413,7 +413,7 @@ uses (for performing 1d FFTs) when running the particle-particle
|
||||
particle-mesh (PPPM) option for long-range Coulombics via the
|
||||
"kspace_style"_kspace_style.html command.
|
||||
|
||||
LAMMPS supports various open-source or vendor-supplied FFT libraries
|
||||
LAMMPS supports common open-source or vendor-supplied FFT libraries
|
||||
for this purpose. If you leave these 3 variables blank, LAMMPS will
|
||||
use the open-source "KISS FFT library"_http://kissfft.sf.net, which is
|
||||
included in the LAMMPS distribution. This library is portable to all
|
||||
@ -423,10 +423,9 @@ package in your build, you can also leave the 3 variables blank.
|
||||
|
||||
Otherwise, select which kinds of FFTs to use as part of the FFT_INC
|
||||
setting by a switch of the form -DFFT_XXX. Recommended values for XXX
|
||||
are: MKL, SCSL, FFTW2, and FFTW3. Legacy options are: INTEL, SGI,
|
||||
ACML, and T3E. For backward compatability, using -DFFT_FFTW will use
|
||||
the FFTW2 library. Using -DFFT_NONE will use the KISS library
|
||||
described above.
|
||||
are: MKL or FFTW3. FFTW2 and NONE are supported as legacy options.
|
||||
Selecting -DFFT_FFTW will use the FFTW3 library and -DFFT_NONE will
|
||||
use the KISS library described above.
|
||||
|
||||
You may also need to set the FFT_INC, FFT_PATH, and FFT_LIB variables,
|
||||
so the compiler and linker can find the needed FFT header and library
|
||||
|
||||
@ -16,10 +16,11 @@ ID, group-ID are documented in "compute"_compute.html command :ulb,l
|
||||
group/group = style name of this compute command :l
|
||||
group2-ID = group ID of second (or same) group :l
|
||||
zero or more keyword/value pairs may be appended :l
|
||||
keyword = {pair} or {kspace} or {boundary} :l
|
||||
keyword = {pair} or {kspace} or {boundary} or {molecule} :l
|
||||
{pair} value = {yes} or {no}
|
||||
{kspace} value = {yes} or {no}
|
||||
{boundary} value = {yes} or {no} :pre
|
||||
{boundary} value = {yes} or {no}
|
||||
{molecule} value = {off} or {inter} or {intra} :pre
|
||||
:ule
|
||||
|
||||
[Examples:]
|
||||
@ -46,6 +47,13 @@ NOTE: The energies computed by the {pair} keyword do not include tail
|
||||
corrections, even if they are enabled via the
|
||||
"pair_modify"_pair_modify.html command.
|
||||
|
||||
If the {molecule} keyword is set to {inter} or {intra} than an
|
||||
additional check is made based on the molecule IDs of the two atoms in
|
||||
each pair before including their pairwise interaction energy and
|
||||
force. For the {inter} setting, the two atoms must be in different
|
||||
molecules. For the {intra} setting, the two atoms must be in the same
|
||||
molecule.
|
||||
|
||||
If the {kspace} keyword is set to {yes}, which is not the default, and
|
||||
if a "kspace_style"_kspace_style.html is defined, then the interaction
|
||||
energy will include a Kspace component which is the long-range
|
||||
@ -66,6 +74,10 @@ affect the force calculation and will be zero if one or both of the
|
||||
groups are charge neutral. This energy correction term is the same as
|
||||
that included in the regular Ewald and PPPM routines.
|
||||
|
||||
NOTE: The {molecule} setting only affects the group/group
|
||||
contributions calculated by the {pair} keyword. It does not affect
|
||||
the group/group contributions calculated by the {kspace} keyword.
|
||||
|
||||
This compute does not calculate any bond or angle or dihedral or
|
||||
improper interactions between atoms in the two groups.
|
||||
|
||||
@ -78,6 +90,22 @@ work (FFTs, Ewald summation) as computing long-range forces for the
|
||||
entire system. Thus it can be costly to invoke this compute too
|
||||
frequently.
|
||||
|
||||
NOTE: If you have a bonded system, then the settings of
|
||||
"special_bonds"_special_bonds.html command can remove pairwise
|
||||
interactions between atoms in the same bond, angle, or dihedral. This
|
||||
is the default setting for the "special_bonds"_special_bonds.html
|
||||
command, and means those pairwise interactions do not appear in the
|
||||
neighbor list. Because this compute uses a neighbor list, it also
|
||||
means those pairs will not be included in the group/group interaction.
|
||||
This does not apply when using long-range coulomb interactions
|
||||
({coul/long}, {coul/msm}, {coul/wolf} or similar. One way to get
|
||||
around this would be to set special_bond scaling factors to very tiny
|
||||
numbers that are not exactly zero (e.g. 1.0e-50). Another workaround
|
||||
is to write a dump file, and use the "rerun"_rerun.html command to
|
||||
compute the group/group interactions for snapshots in the dump file.
|
||||
The rerun script can use a "special_bonds"_special_bonds.html command
|
||||
that includes all pairs in the neighbor list.
|
||||
|
||||
If you desire a breakdown of the interactions into a pairwise and
|
||||
Kspace component, simply invoke the compute twice with the appropriate
|
||||
yes/no settings for the {pair} and {kspace} keywords. This is no more
|
||||
@ -119,7 +147,8 @@ The {ewald} and {pppm} styles do.
|
||||
|
||||
[Default:]
|
||||
|
||||
The option defaults are pair = yes, kspace = no, and boundary = yes.
|
||||
The option defaults are pair = yes, kspace = no, boundary = yes,
|
||||
molecule = off.
|
||||
|
||||
:line
|
||||
|
||||
|
||||
@ -41,14 +41,14 @@ NOTE: If you have a bonded system, then the settings of
|
||||
interactions between atoms in the same bond, angle, or dihedral. This
|
||||
is the default setting for the "special_bonds"_special_bonds.html
|
||||
command, and means those pairwise interactions do not appear in the
|
||||
neighbor list. Because this fix uses the neighbor list, it also means
|
||||
neighbor list. Because this fix uses a neighbor list, it also means
|
||||
those pairs will not be included in the RDF. This does not apply when
|
||||
using long-range coulomb ({coul/long}, {coul/msm}, {coul/wolf} or
|
||||
similar. One way to get around this would be to set special_bond
|
||||
scaling factors to very tiny numbers that are not exactly zero
|
||||
(e.g. 1.0e-50). Another workaround is to write a dump file, and use
|
||||
the "rerun"_rerun.html command to compute the RDF for snapshots in the
|
||||
dump file. The rerun script can use a
|
||||
using long-range coulomb interactions ({coul/long}, {coul/msm},
|
||||
{coul/wolf} or similar. One way to get around this would be to set
|
||||
special_bond scaling factors to very tiny numbers that are not exactly
|
||||
zero (e.g. 1.0e-50). Another workaround is to write a dump file, and
|
||||
use the "rerun"_rerun.html command to compute the RDF for snapshots in
|
||||
the dump file. The rerun script can use a
|
||||
"special_bonds"_special_bonds.html command that includes all pairs in
|
||||
the neighbor list.
|
||||
|
||||
|
||||
@ -13,6 +13,7 @@ pair_style morse/opt command :h3
|
||||
pair_style morse/smooth/linear command :h3
|
||||
pair_style morse/smooth/linear/omp command :h3
|
||||
pair_style morse/soft command :h3
|
||||
pair_style morse/kk command :h3
|
||||
|
||||
[Syntax:]
|
||||
|
||||
|
||||
@ -54,7 +54,8 @@ reset_timestep 0
|
||||
variable pxy equal pxy
|
||||
variable pxx equal pxx-press
|
||||
|
||||
fix avstress all ave/time $s $p $d v_pxy v_pxx ave one file einstein.dat
|
||||
fix avstress all ave/time $s $p $d v_pxy v_pxx ave one &
|
||||
file profile.einstein.2d
|
||||
|
||||
# Diagonal components of SS are larger by factor 2-2/d,
|
||||
# which is 4/3 for d=3, but 1 for d=2.
|
||||
|
||||
189
examples/coreshell/log.9Nov16.coreshell.dsf.g++.4
Normal file
189
examples/coreshell/log.9Nov16.coreshell.dsf.g++.4
Normal file
@ -0,0 +1,189 @@
|
||||
LAMMPS (26 Jan 2017)
|
||||
# Testsystem for core-shell model compared to Mitchel and Finchham
|
||||
# Hendrik Heenen, June 2014
|
||||
|
||||
# ------------------------ INITIALIZATION ----------------------------
|
||||
|
||||
units metal
|
||||
dimension 3
|
||||
boundary p p p
|
||||
atom_style full
|
||||
|
||||
# ----------------------- ATOM DEFINITION ----------------------------
|
||||
|
||||
fix csinfo all property/atom i_CSID
|
||||
read_data data.coreshell fix csinfo NULL CS-Info
|
||||
orthogonal box = (0 0 0) to (24.096 24.096 24.096)
|
||||
1 by 2 by 2 MPI processor grid
|
||||
reading atoms ...
|
||||
432 atoms
|
||||
scanning bonds ...
|
||||
1 = max bonds/atom
|
||||
reading bonds ...
|
||||
216 bonds
|
||||
1 = max # of 1-2 neighbors
|
||||
0 = max # of 1-3 neighbors
|
||||
0 = max # of 1-4 neighbors
|
||||
1 = max # of special neighbors
|
||||
|
||||
group cores type 1 2
|
||||
216 atoms in group cores
|
||||
group shells type 3 4
|
||||
216 atoms in group shells
|
||||
|
||||
neighbor 2.0 bin
|
||||
comm_modify vel yes
|
||||
|
||||
# ------------------------ FORCE FIELDS ------------------------------
|
||||
|
||||
pair_style born/coul/dsf/cs 0.1 20.0 20.0 # A, rho, sigma=0, C, D
|
||||
pair_coeff * * 0.0 1.000 0.00 0.00 0.00
|
||||
pair_coeff 3 3 487.0 0.23768 0.00 1.05 0.50 #Na-Na
|
||||
pair_coeff 3 4 145134.0 0.23768 0.00 6.99 8.70 #Na-Cl
|
||||
pair_coeff 4 4 405774.0 0.23768 0.00 72.40 145.40 #Cl-Cl
|
||||
|
||||
bond_style harmonic
|
||||
bond_coeff 1 63.014 0.0
|
||||
bond_coeff 2 25.724 0.0
|
||||
|
||||
# ------------------------ Equilibration Run -------------------------------
|
||||
|
||||
reset_timestep 0
|
||||
|
||||
thermo 50
|
||||
thermo_style custom step etotal pe ke temp press epair evdwl ecoul elong ebond fnorm fmax vol
|
||||
|
||||
compute CSequ all temp/cs cores shells
|
||||
|
||||
# output via chunk method
|
||||
|
||||
#compute prop all property/atom i_CSID
|
||||
#compute cs_chunk all chunk/atom c_prop
|
||||
#compute cstherm all temp/chunk cs_chunk temp internal com yes cdof 3.0
|
||||
#fix ave_chunk all ave/time 100 1 100 c_cstherm file chunk.dump mode vector
|
||||
|
||||
thermo_modify temp CSequ
|
||||
|
||||
# velocity bias option
|
||||
|
||||
velocity all create 1427 134 dist gaussian mom yes rot no bias yes temp CSequ
|
||||
Neighbor list info ...
|
||||
update every 1 steps, delay 10 steps, check yes
|
||||
max neighbors/atom: 2000, page size: 100000
|
||||
master list distance cutoff = 22
|
||||
ghost atom cutoff = 22
|
||||
binsize = 11, bins = 3 3 3
|
||||
1 neighbor lists, perpetual/occasional/extra = 1 0 0
|
||||
(1) pair born/coul/dsf/cs, half, perpetual
|
||||
pair build: half/bin/newton
|
||||
stencil: half/bin/3d/newton
|
||||
bin: standard
|
||||
velocity all scale 1427 temp CSequ
|
||||
|
||||
fix thermoberendsen all temp/berendsen 1427 1427 0.4
|
||||
fix nve all nve
|
||||
fix_modify thermoberendsen temp CSequ
|
||||
|
||||
# 2 fmsec timestep
|
||||
|
||||
timestep 0.002
|
||||
run 500
|
||||
Memory usage per processor = 6.8559 Mbytes
|
||||
Step TotEng PotEng KinEng Temp Press E_pair E_vdwl E_coul E_long E_bond Fnorm Fmax Volume
|
||||
0 -635.80596 -675.46362 39.657659 1427 -21302.622 -675.46362 1.6320365 -677.09565 0 0 1.5814015e-14 3.2317898e-15 13990.5
|
||||
50 -634.07021 -666.11867 32.048452 1153.1982 -4560.945 -668.28236 37.756542 -706.0389 0 2.163691 13.802484 3.022372 13990.5
|
||||
100 -631.97128 -662.02544 30.054164 1081.4378 -3497.564 -664.61825 39.275003 -703.89325 0 2.5928078 13.956833 2.5417699 13990.5
|
||||
150 -630.14953 -663.04215 32.892622 1183.5739 -88.43828 -665.63444 46.239965 -711.87441 0 2.5922927 14.667898 2.4964255 13990.5
|
||||
200 -628.52878 -663.9795 35.45072 1275.6219 -1755.9004 -666.73564 41.758052 -708.49369 0 2.7561421 14.230743 3.0924004 13990.5
|
||||
250 -627.27102 -662.025 34.753978 1250.5511 -1234.0918 -665.13519 43.170874 -708.30606 0 3.1101887 14.221086 1.941354 13990.5
|
||||
300 -626.5495 -663.74287 37.193368 1338.3275 -2049.3444 -666.45574 40.476148 -706.93188 0 2.7128711 13.330425 1.7756755 13990.5
|
||||
350 -625.87313 -665.21855 39.345421 1415.7647 -1543.1723 -667.90872 41.577366 -709.48609 0 2.6901682 13.541311 1.854662 13990.5
|
||||
400 -625.09344 -661.26404 36.1706 1301.5253 -729.96729 -664.10334 43.468765 -707.57211 0 2.8392963 13.663555 1.9067551 13990.5
|
||||
450 -624.46214 -660.01362 35.551477 1279.2474 -1617.7158 -663.06571 41.644856 -704.71057 0 3.0520921 14.527005 1.7280213 13990.5
|
||||
500 -623.49246 -659.2527 35.76024 1286.7593 -935.99238 -662.32953 43.038808 -705.36834 0 3.0768302 14.099593 1.9831106 13990.5
|
||||
Loop time of 4.09864 on 4 procs for 500 steps with 432 atoms
|
||||
|
||||
Performance: 21.080 ns/day, 1.139 hours/ns, 121.992 timesteps/s
|
||||
99.7% CPU use with 4 MPI tasks x no OpenMP threads
|
||||
|
||||
MPI task timing breakdown:
|
||||
Section | min time | avg time | max time |%varavg| %total
|
||||
---------------------------------------------------------------
|
||||
Pair | 3.3804 | 3.568 | 3.8354 | 8.9 | 87.05
|
||||
Bond | 0.00074339 | 0.00079519 | 0.00087976 | 0.0 | 0.02
|
||||
Neigh | 0.045851 | 0.046084 | 0.046361 | 0.1 | 1.12
|
||||
Comm | 0.20413 | 0.47123 | 0.65875 | 24.3 | 11.50
|
||||
Output | 0.00044298 | 0.00046057 | 0.00051165 | 0.0 | 0.01
|
||||
Modify | 0.0064909 | 0.0067219 | 0.0069766 | 0.2 | 0.16
|
||||
Other | | 0.005345 | | | 0.13
|
||||
|
||||
Nlocal: 108 ave 114 max 105 min
|
||||
Histogram: 1 1 1 0 0 0 0 0 0 1
|
||||
Nghost: 6527 ave 6599 max 6472 min
|
||||
Histogram: 1 0 1 0 1 0 0 0 0 1
|
||||
Neighs: 74388.2 ave 75855 max 73680 min
|
||||
Histogram: 1 2 0 0 0 0 0 0 0 1
|
||||
|
||||
Total # of neighbors = 297553
|
||||
Ave neighs/atom = 688.78
|
||||
Ave special neighs/atom = 1
|
||||
Neighbor list builds = 20
|
||||
Dangerous builds = 0
|
||||
|
||||
unfix thermoberendsen
|
||||
|
||||
# ------------------------ Dynamic Run -------------------------------
|
||||
|
||||
run 1000
|
||||
Memory usage per processor = 6.85787 Mbytes
|
||||
Step TotEng PotEng KinEng Temp Press E_pair E_vdwl E_coul E_long E_bond Fnorm Fmax Volume
|
||||
500 -623.49319 -659.2527 35.759511 1286.7331 -936.04802 -662.32953 43.038808 -705.36834 0 3.0768302 14.099593 1.9831106 13990.5
|
||||
550 -623.44059 -663.57938 40.138795 1444.3127 -935.73484 -666.2789 42.563337 -708.84224 0 2.6995167 13.918509 2.3189805 13990.5
|
||||
600 -623.4703 -660.01592 36.545618 1315.0196 1327.3492 -663.08845 47.985462 -711.07391 0 3.0725254 15.192713 2.4098428 13990.5
|
||||
650 -623.46796 -661.56776 38.099807 1370.9439 457.82439 -664.81976 45.495622 -710.31538 0 3.2519966 15.026057 1.8500226 13990.5
|
||||
700 -623.50158 -659.5131 36.011523 1295.8012 -460.03772 -663.1078 43.938203 -707.046 0 3.5946908 14.660979 2.4825518 13990.5
|
||||
750 -623.44787 -661.93353 38.485658 1384.8279 97.429626 -664.9551 45.083146 -710.03825 0 3.0215753 15.10043 2.3433897 13990.5
|
||||
800 -623.48215 -659.50655 36.024402 1296.2647 1097.3866 -662.61124 47.251998 -709.86324 0 3.1046914 14.556382 2.0543766 13990.5
|
||||
850 -623.45868 -661.13782 37.679134 1355.8068 -1802.1624 -664.41257 40.70845 -705.12102 0 3.2747525 14.691444 2.2054332 13990.5
|
||||
900 -623.43556 -663.59137 40.155815 1444.9251 534.99197 -666.71877 45.601619 -712.32039 0 3.127395 14.741411 2.5807895 13990.5
|
||||
950 -623.51318 -661.57916 38.06598 1369.7267 -678.12625 -664.37535 43.207862 -707.58322 0 2.7961988 14.430307 2.3936105 13990.5
|
||||
1000 -623.47287 -661.22274 37.749874 1358.3523 634.7979 -664.42973 46.373361 -710.80309 0 3.2069879 15.891192 2.4042765 13990.5
|
||||
1050 -623.48133 -661.52868 38.047347 1369.0562 -583.15228 -664.6098 43.618772 -708.22857 0 3.081116 14.806856 2.3447613 13990.5
|
||||
1100 -623.47867 -661.83761 38.358946 1380.2685 -868.9779 -664.8826 42.84846 -707.73106 0 3.044983 14.69567 2.399143 13990.5
|
||||
1150 -623.44713 -661.21299 37.765857 1358.9274 405.14554 -664.09567 45.578739 -709.6744 0 2.8826753 15.437367 3.1381305 13990.5
|
||||
1200 -623.46549 -660.91706 37.451568 1347.6183 699.78996 -664.0883 46.36297 -710.45127 0 3.1712473 15.109665 1.8891886 13990.5
|
||||
1250 -623.49296 -658.2218 34.728838 1249.6464 1061.0154 -661.29052 47.668699 -708.95922 0 3.0687228 14.901367 2.3964137 13990.5
|
||||
1300 -623.49837 -660.91022 37.411844 1346.1889 226.99512 -664.35989 45.352287 -709.71217 0 3.4496704 15.161542 2.2137993 13990.5
|
||||
1350 -623.46718 -658.80365 35.336469 1271.5108 1039.6469 -662.16908 47.565671 -709.73475 0 3.3654314 15.892516 2.7888426 13990.5
|
||||
1400 -623.47124 -661.45375 37.982513 1366.7233 -379.56023 -664.6321 43.788306 -708.42041 0 3.1783497 14.251126 1.7415409 13990.5
|
||||
1450 -623.46671 -660.17518 36.708464 1320.8792 -374.37056 -662.92706 44.083648 -707.01071 0 2.7518803 15.210167 1.9984277 13990.5
|
||||
1500 -623.50515 -659.06488 35.559725 1279.5442 260.37822 -662.39548 45.779764 -708.17524 0 3.3306005 14.682396 2.4201107 13990.5
|
||||
Loop time of 8.26746 on 4 procs for 1000 steps with 432 atoms
|
||||
|
||||
Performance: 20.901 ns/day, 1.148 hours/ns, 120.956 timesteps/s
|
||||
99.7% CPU use with 4 MPI tasks x no OpenMP threads
|
||||
|
||||
MPI task timing breakdown:
|
||||
Section | min time | avg time | max time |%varavg| %total
|
||||
---------------------------------------------------------------
|
||||
Pair | 6.706 | 7.1568 | 7.6597 | 12.7 | 86.57
|
||||
Bond | 0.0014617 | 0.0015531 | 0.0016506 | 0.2 | 0.02
|
||||
Neigh | 0.10511 | 0.10522 | 0.10532 | 0.0 | 1.27
|
||||
Comm | 0.48547 | 0.98841 | 1.4393 | 34.0 | 11.96
|
||||
Output | 0.0012085 | 0.0012462 | 0.0013196 | 0.1 | 0.02
|
||||
Modify | 0.0021446 | 0.0021989 | 0.0022545 | 0.1 | 0.03
|
||||
Other | | 0.01204 | | | 0.15
|
||||
|
||||
Nlocal: 108 ave 114 max 94 min
|
||||
Histogram: 1 0 0 0 0 0 0 0 1 2
|
||||
Nghost: 6512.25 ave 6586 max 6456 min
|
||||
Histogram: 1 0 0 2 0 0 0 0 0 1
|
||||
Neighs: 74248.2 ave 77441 max 65858 min
|
||||
Histogram: 1 0 0 0 0 0 0 0 0 3
|
||||
|
||||
Total # of neighbors = 296993
|
||||
Ave neighs/atom = 687.484
|
||||
Ave special neighs/atom = 1
|
||||
Neighbor list builds = 46
|
||||
Dangerous builds = 0
|
||||
Total wall time: 0:00:12
|
||||
17
potentials/Ge.tersoff
Normal file
17
potentials/Ge.tersoff
Normal file
@ -0,0 +1,17 @@
|
||||
# DATE: 2016-12-21 CONTRIBUTOR: Sayyed Jalil Mahdizadh, saja.mahdizadeh@gmail.com CITATION: Mahdizadeh, Akhlamadi, J. Mol. Graph. Model. 72, 1-5 (2017)
|
||||
|
||||
# Tersoff parameters for various elements and mixtures
|
||||
# multiple entries can be added to this file, LAMMPS reads the ones it needs
|
||||
# these entries are in LAMMPS "metal" units:
|
||||
# A,B = eV; lambda1,lambda2,lambda3 = 1/Angstroms; R,D = Angstroms
|
||||
# other quantities are unitless
|
||||
|
||||
# This is the Ge reparameterization from the following paper:
|
||||
# Optimized Tersoff empirical potential for germanene
|
||||
# Mahdizadeh, Akhlamadi, J. Mol. Graph. Model. 72, 1-5 (2017)
|
||||
|
||||
# format of a single entry (one or more lines):
|
||||
# element 1, element 2, element 3,
|
||||
# m, gamma, lambda3, c, d, costheta0, n, beta, lambda2, B, R, D, lambda1, A
|
||||
|
||||
Ge Ge Ge 3.0 1.0 0.0 1.0643e5 15.2 -0.35 0.75627 5.017e-7 1.71 430.0 2.95 0.15 2.4451 1760.1
|
||||
@ -278,7 +278,7 @@ void FixWallGranRegion::update_contacts(int i, int nc)
|
||||
iold = 0;
|
||||
while (iold < ncontact[i]) {
|
||||
for (m = 0; m < nc; m++)
|
||||
if (region->contact[m].iwall = walls[i][iold]) break;
|
||||
if (region->contact[m].iwall == walls[i][iold]) break;
|
||||
if (m < nc) {
|
||||
ilast = ncontact[i]-1;
|
||||
for (j = 0; j < sheardim; j++)
|
||||
|
||||
@ -799,3 +799,14 @@ double PairGranHookeHistory::memory_usage()
|
||||
double bytes = nmax * sizeof(double);
|
||||
return bytes;
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
return ptr to FixShearHistory class
|
||||
called by Neighbor when setting up neighbor lists
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
void *PairGranHookeHistory::extract(const char *str, int &dim)
|
||||
{
|
||||
if (strcmp(str,"history") == 0) return (void *) fix_history;
|
||||
return NULL;
|
||||
}
|
||||
|
||||
@ -43,6 +43,7 @@ class PairGranHookeHistory : public Pair {
|
||||
int pack_forward_comm(int, int *, double *, int, int *);
|
||||
void unpack_forward_comm(int, int, double *);
|
||||
double memory_usage();
|
||||
void *extract(const char *, int &);
|
||||
|
||||
protected:
|
||||
double kn,kt,gamman,gammat,xmu;
|
||||
|
||||
@ -173,6 +173,8 @@ action pair_lj_gromacs_kokkos.cpp
|
||||
action pair_lj_gromacs_kokkos.h
|
||||
action pair_lj_sdk_kokkos.cpp pair_lj_sdk.cpp
|
||||
action pair_lj_sdk_kokkos.h pair_lj_sdk.h
|
||||
action pair_morse_kokkos.cpp
|
||||
action pair_morse_kokkos.h
|
||||
action pair_reax_c_kokkos.cpp pair_reax_c.cpp
|
||||
action pair_reax_c_kokkos.h pair_reax_c.h
|
||||
action pair_sw_kokkos.cpp pair_sw.cpp
|
||||
|
||||
@ -70,23 +70,20 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
if(k_eatom.dimension_0()<maxeatom) {
|
||||
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
}
|
||||
d_eatom = k_eatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
if (vflag_atom) {
|
||||
if(k_vatom.dimension_0()<maxvatom) {
|
||||
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
}
|
||||
d_vatom = k_vatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
|
||||
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
else atomKK->modified(execution_space,F_MASK);
|
||||
|
||||
x = atomKK->k_x.view<DeviceType>();
|
||||
f = atomKK->k_f.view<DeviceType>();
|
||||
neighborKK->k_anglelist.template sync<DeviceType>();
|
||||
|
||||
@ -957,10 +957,10 @@ struct AtomVecAngleKokkos_UnpackBorder {
|
||||
_x(i+_first,0) = _buf(i,0);
|
||||
_x(i+_first,1) = _buf(i,1);
|
||||
_x(i+_first,2) = _buf(i,2);
|
||||
_tag(i+_first) = static_cast<int> (_buf(i,3));
|
||||
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
|
||||
_type(i+_first) = static_cast<int> (_buf(i,4));
|
||||
_mask(i+_first) = static_cast<int> (_buf(i,5));
|
||||
_molecule(i+_first) = static_cast<int> (_buf(i,6));
|
||||
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
|
||||
|
||||
}
|
||||
};
|
||||
|
||||
@ -836,7 +836,7 @@ struct AtomVecAtomicKokkos_UnpackBorder {
|
||||
_x(i+_first,0) = _buf(i,0);
|
||||
_x(i+_first,1) = _buf(i,1);
|
||||
_x(i+_first,2) = _buf(i,2);
|
||||
_tag(i+_first) = static_cast<int> (_buf(i,3));
|
||||
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
|
||||
_type(i+_first) = static_cast<int> (_buf(i,4));
|
||||
_mask(i+_first) = static_cast<int> (_buf(i,5));
|
||||
// printf("%i %i %lf %lf %lf %i BORDER\n",_tag(i+_first),i+_first,_x(i+_first,0),_x(i+_first,1),_x(i+_first,2),_type(i+_first));
|
||||
|
||||
@ -905,10 +905,10 @@ struct AtomVecBondKokkos_UnpackBorder {
|
||||
_x(i+_first,0) = _buf(i,0);
|
||||
_x(i+_first,1) = _buf(i,1);
|
||||
_x(i+_first,2) = _buf(i,2);
|
||||
_tag(i+_first) = static_cast<int> (_buf(i,3));
|
||||
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
|
||||
_type(i+_first) = static_cast<int> (_buf(i,4));
|
||||
_mask(i+_first) = static_cast<int> (_buf(i,5));
|
||||
_molecule(i+_first) = static_cast<int> (_buf(i,6));
|
||||
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
|
||||
|
||||
}
|
||||
};
|
||||
|
||||
@ -872,7 +872,7 @@ struct AtomVecChargeKokkos_UnpackBorder {
|
||||
_x(i+_first,0) = _buf(i,0);
|
||||
_x(i+_first,1) = _buf(i,1);
|
||||
_x(i+_first,2) = _buf(i,2);
|
||||
_tag(i+_first) = static_cast<int> (_buf(i,3));
|
||||
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
|
||||
_type(i+_first) = static_cast<int> (_buf(i,4));
|
||||
_mask(i+_first) = static_cast<int> (_buf(i,5));
|
||||
_q(i+_first) = _buf(i,6);
|
||||
|
||||
@ -1029,10 +1029,10 @@ struct AtomVecMolecularKokkos_UnpackBorder {
|
||||
_x(i+_first,0) = _buf(i,0);
|
||||
_x(i+_first,1) = _buf(i,1);
|
||||
_x(i+_first,2) = _buf(i,2);
|
||||
_tag(i+_first) = static_cast<int> (_buf(i,3));
|
||||
_tag(i+_first) = static_cast<tagint> (_buf(i,3));
|
||||
_type(i+_first) = static_cast<int> (_buf(i,4));
|
||||
_mask(i+_first) = static_cast<int> (_buf(i,5));
|
||||
_molecule(i+_first) = static_cast<int> (_buf(i,6));
|
||||
_molecule(i+_first) = static_cast<tagint> (_buf(i,6));
|
||||
|
||||
}
|
||||
};
|
||||
|
||||
@ -77,21 +77,18 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"bond:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"bond:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_k.template sync<DeviceType>();
|
||||
k_r0.template sync<DeviceType>();
|
||||
k_epsilon.template sync<DeviceType>();
|
||||
k_sigma.template sync<DeviceType>();
|
||||
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
else atomKK->modified(execution_space,F_MASK);
|
||||
|
||||
x = atomKK->k_x.view<DeviceType>();
|
||||
f = atomKK->k_f.view<DeviceType>();
|
||||
|
||||
@ -66,8 +66,8 @@ class BondFENEKokkos : public BondFENE {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
DAT::tdual_int_scalar k_warning_flag;
|
||||
typename AT::t_int_scalar d_warning_flag;
|
||||
|
||||
@ -67,18 +67,18 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
if(k_eatom.dimension_0()<maxeatom) {
|
||||
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
}
|
||||
d_eatom = k_eatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
if (vflag_atom) {
|
||||
if(k_vatom.dimension_0()<maxvatom) {
|
||||
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
}
|
||||
d_vatom = k_vatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
|
||||
// if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
|
||||
@ -80,29 +80,24 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
if(k_eatom.dimension_0()<maxeatom) {
|
||||
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.template view<DeviceType>();
|
||||
k_eatom_pair = Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,DeviceType>("dihedral:eatom_pair",maxeatom);
|
||||
d_eatom_pair = k_eatom.d_view;
|
||||
}
|
||||
d_eatom_pair = k_eatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
if (vflag_atom) {
|
||||
if(k_vatom.dimension_0()<maxvatom) {
|
||||
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.template view<DeviceType>();
|
||||
k_vatom_pair = Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,DeviceType>("dihedral:vatom_pair",maxvatom);
|
||||
d_vatom_pair = k_vatom.d_view;
|
||||
}
|
||||
d_vatom_pair = k_vatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
|
||||
|
||||
//atomKK->sync(execution_space,datamask_read);
|
||||
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
else atomKK->modified(execution_space,F_MASK);
|
||||
|
||||
x = atomKK->k_x.view<DeviceType>();
|
||||
f = atomKK->k_f.view<DeviceType>();
|
||||
q = atomKK->k_q.view<DeviceType>();
|
||||
|
||||
@ -77,21 +77,18 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_k1.template sync<DeviceType>();
|
||||
k_k2.template sync<DeviceType>();
|
||||
k_k3.template sync<DeviceType>();
|
||||
k_k4.template sync<DeviceType>();
|
||||
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
else atomKK->modified(execution_space,F_MASK);
|
||||
|
||||
x = atomKK->k_x.view<DeviceType>();
|
||||
f = atomKK->k_f.view<DeviceType>();
|
||||
|
||||
@ -68,8 +68,8 @@ class DihedralOPLSKokkos : public DihedralOPLS {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
int nlocal,newton_bond;
|
||||
int eflag,vflag;
|
||||
|
||||
@ -250,7 +250,7 @@ struct DomainPBCFunctor {
|
||||
x(i,0) += period[0];
|
||||
if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) += h_rate[0];
|
||||
imageint idim = image[i] & IMGMASK;
|
||||
const int otherdims = image[i] ^ idim;
|
||||
const imageint otherdims = image[i] ^ idim;
|
||||
idim--;
|
||||
idim &= IMGMASK;
|
||||
image[i] = otherdims | idim;
|
||||
@ -260,7 +260,7 @@ struct DomainPBCFunctor {
|
||||
x(i,0) = MAX(x(i,0),lo[0]);
|
||||
if (DEFORM_VREMAP && (mask[i] & deform_groupbit)) v(i,0) -= h_rate[0];
|
||||
imageint idim = image[i] & IMGMASK;
|
||||
const int otherdims = image[i] ^ idim;
|
||||
const imageint otherdims = image[i] ^ idim;
|
||||
idim++;
|
||||
idim &= IMGMASK;
|
||||
image[i] = otherdims | idim;
|
||||
|
||||
@ -387,7 +387,7 @@ KOKKOS_INLINE_FUNCTION
|
||||
void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const bool &final) const
|
||||
{
|
||||
const int i = d_ilist[ii];
|
||||
int j,jj,jtag,jtype,flag;
|
||||
int j,jj,jtype,flag;
|
||||
|
||||
if (mask[i] & groupbit) {
|
||||
|
||||
@ -395,7 +395,7 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int jnum = d_numneigh[i];
|
||||
if (final)
|
||||
d_firstnbr[i] = m_fill;
|
||||
@ -403,7 +403,6 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
|
||||
for (jj = 0; jj < jnum; jj++) {
|
||||
j = d_neighbors(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
|
||||
jtype = type(j);
|
||||
|
||||
const X_FLOAT delx = x(j,0) - xtmp;
|
||||
@ -411,10 +410,11 @@ void FixQEqReaxKokkos<DeviceType>::compute_h_item(int ii, int &m_fill, const boo
|
||||
const X_FLOAT delz = x(j,2) - ztmp;
|
||||
|
||||
if (neighflag != FULL) {
|
||||
const tagint jtag = tag(j);
|
||||
flag = 0;
|
||||
if (j < nlocal) flag = 1;
|
||||
else if (tag[i] < tag[j]) flag = 1;
|
||||
else if (tag[i] == tag[j]) {
|
||||
else if (itag < jtag) flag = 1;
|
||||
else if (itag == jtag) {
|
||||
if (delz > SMALL) flag = 1;
|
||||
else if (fabs(delz) < SMALL) {
|
||||
if (dely > SMALL) flag = 1;
|
||||
|
||||
@ -159,7 +159,8 @@ class FixQEqReaxKokkos : public FixQEqReax {
|
||||
//typename ArrayTypes<DeviceType>::t_float_1d_randomread mass, q;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread mass;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d q;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d type, tag, mask;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d type, mask;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
|
||||
DAT::tdual_float_1d k_q;
|
||||
typename AT::t_float_1d d_q;
|
||||
|
||||
@ -76,7 +76,6 @@ void FixReaxCBondsKokkos::init()
|
||||
void FixReaxCBondsKokkos::Output_ReaxC_Bonds(bigint ntimestep, FILE *fp)
|
||||
|
||||
{
|
||||
int i, j;
|
||||
int nbuf_local;
|
||||
int nlocal_max, numbonds, numbonds_max;
|
||||
double *buf;
|
||||
|
||||
@ -77,18 +77,18 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
if(k_eatom.dimension_0()<maxeatom) {
|
||||
//if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
}
|
||||
d_eatom = k_eatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
if (vflag_atom) {
|
||||
if(k_vatom.dimension_0()<maxvatom) {
|
||||
//if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
}
|
||||
d_vatom = k_vatom.template view<DeviceType>();
|
||||
//}
|
||||
}
|
||||
|
||||
//atomKK->sync(execution_space,datamask_read);
|
||||
|
||||
@ -24,8 +24,8 @@ namespace LAMMPS_NS {
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::NPairKokkos(LAMMPS *lmp) : NPair(lmp) {
|
||||
|
||||
}
|
||||
|
||||
@ -33,8 +33,8 @@ NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::NPairKokkos(LAMMPS *lmp) : NPair(lmp)
|
||||
copy needed info from Neighbor class to this build class
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_neighbor_info()
|
||||
{
|
||||
NPair::copy_neighbor_info();
|
||||
|
||||
@ -62,8 +62,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_neighbor_info()
|
||||
copy per-atom and per-bin vectors from NBin class to this build class
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_bin_info()
|
||||
{
|
||||
NPair::copy_bin_info();
|
||||
|
||||
@ -78,8 +78,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_bin_info()
|
||||
copy needed info from NStencil class to this build class
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::copy_stencil_info()
|
||||
{
|
||||
NPair::copy_stencil_info();
|
||||
|
||||
@ -106,8 +106,8 @@ void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::copy_stencil_info()
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST>::build(NeighList *list_)
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
void NPairKokkos<DeviceType,HALF_NEIGH,GHOST,TRI>::build(NeighList *list_)
|
||||
{
|
||||
NeighListKokkos<DeviceType>* list = (NeighListKokkos<DeviceType>*) list_;
|
||||
const int nlocal = includegroup?atom->nfirst:atom->nlocal;
|
||||
@ -196,14 +196,14 @@ if (GHOST) {
|
||||
Kokkos::parallel_for(nall, f);
|
||||
} else {
|
||||
if (newton_pair) {
|
||||
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,1> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
|
||||
NPairKokkosBuildFunctor<DeviceType,TRI?0:HALF_NEIGH,1,TRI> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
Kokkos::parallel_for(config, f);
|
||||
#else
|
||||
Kokkos::parallel_for(nall, f);
|
||||
#endif
|
||||
} else {
|
||||
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
|
||||
NPairKokkosBuildFunctor<DeviceType,HALF_NEIGH,0,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
Kokkos::parallel_for(config, f);
|
||||
#else
|
||||
@ -293,7 +293,7 @@ int NeighborKokkosExecute<DeviceType>::exclusion(const int &i,const int &j,
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType> template<int HalfNeigh,int Newton>
|
||||
template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>
|
||||
void NeighborKokkosExecute<DeviceType>::
|
||||
build_Item(const int &i) const
|
||||
{
|
||||
@ -365,7 +365,7 @@ void NeighborKokkosExecute<DeviceType>::
|
||||
const int jbin = ibin + stencil[k];
|
||||
|
||||
// get subview of jbin
|
||||
if(HalfNeigh&&(ibin==jbin)) continue;
|
||||
if(HalfNeigh && (ibin==jbin)) continue;
|
||||
//const ArrayTypes<DeviceType>::t_int_1d_const_um =Kokkos::subview<t_int_1d_const_um>(bins,jbin,ALL);
|
||||
for(int m = 0; m < c_bincount(jbin); m++) {
|
||||
|
||||
@ -374,6 +374,16 @@ void NeighborKokkosExecute<DeviceType>::
|
||||
|
||||
if(HalfNeigh && !Newton && (j < i)) continue;
|
||||
if(!HalfNeigh && j==i) continue;
|
||||
if(Tri) {
|
||||
if (x(j,2) < ztmp) continue;
|
||||
if (x(j,2) == ztmp) {
|
||||
if (x(j,1) < ytmp) continue;
|
||||
if (x(j,1) == ytmp) {
|
||||
if (x(j,0) < xtmp) continue;
|
||||
if (x(j,0) == xtmp && j <= i) continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
if(exclude && exclusion(i,j,itype,jtype)) continue;
|
||||
|
||||
const X_FLOAT delx = xtmp - x(j, 0);
|
||||
@ -428,7 +438,7 @@ extern __shared__ X_FLOAT sharedmem[];
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType> template<int HalfNeigh,int Newton>
|
||||
template<class DeviceType> template<int HalfNeigh,int Newton,int Tri>
|
||||
__device__ inline
|
||||
void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const
|
||||
{
|
||||
@ -491,6 +501,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli
|
||||
((j >= nlocal) && ((x(j, 2) < ztmp) || (x(j, 2) == ztmp && x(j, 1) < ytmp) ||
|
||||
(x(j, 2) == ztmp && x(j, 1) == ytmp && x(j, 0) < xtmp)))))
|
||||
) continue;
|
||||
if(Tri) {
|
||||
if (x(j,2) < ztmp) continue;
|
||||
if (x(j,2) == ztmp) {
|
||||
if (x(j,1) < ytmp) continue;
|
||||
if (x(j,1) == ytmp) {
|
||||
if (x(j,0) < xtmp) continue;
|
||||
if (x(j,0) == xtmp && j <= i) continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
if(exclude && exclusion(i,j,itype,jtype)) continue;
|
||||
const X_FLOAT delx = xtmp - other_x[m];
|
||||
const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];
|
||||
@ -558,6 +578,16 @@ void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPoli
|
||||
//if(HalfNeigh && (j < i)) continue;
|
||||
if(HalfNeigh && !Newton && (j < i)) continue;
|
||||
if(!HalfNeigh && j==i) continue;
|
||||
if(Tri) {
|
||||
if (x(j,2) < ztmp) continue;
|
||||
if (x(j,2) == ztmp) {
|
||||
if (x(j,1) < ytmp) continue;
|
||||
if (x(j,1) == ytmp) {
|
||||
if (x(j,0) < xtmp) continue;
|
||||
if (x(j,0) == xtmp && j <= i) continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
if(exclude && exclusion(i,j,itype,jtype)) continue;
|
||||
|
||||
const X_FLOAT delx = xtmp - other_x[m];
|
||||
@ -736,14 +766,16 @@ void NeighborKokkosExecute<DeviceType>::
|
||||
}
|
||||
|
||||
namespace LAMMPS_NS {
|
||||
template class NPairKokkos<LMPDeviceType,0,0>;
|
||||
template class NPairKokkos<LMPDeviceType,0,1>;
|
||||
template class NPairKokkos<LMPDeviceType,1,0>;
|
||||
template class NPairKokkos<LMPDeviceType,1,1>;
|
||||
template class NPairKokkos<LMPDeviceType,0,0,0>;
|
||||
template class NPairKokkos<LMPDeviceType,0,1,0>;
|
||||
template class NPairKokkos<LMPDeviceType,1,0,0>;
|
||||
template class NPairKokkos<LMPDeviceType,1,1,0>;
|
||||
template class NPairKokkos<LMPDeviceType,1,0,1>;
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
template class NPairKokkos<LMPHostType,0,0>;
|
||||
template class NPairKokkos<LMPHostType,0,1>;
|
||||
template class NPairKokkos<LMPHostType,1,0>;
|
||||
template class NPairKokkos<LMPHostType,1,1>;
|
||||
template class NPairKokkos<LMPHostType,0,0,0>;
|
||||
template class NPairKokkos<LMPHostType,0,1,0>;
|
||||
template class NPairKokkos<LMPHostType,1,0,0>;
|
||||
template class NPairKokkos<LMPHostType,1,1,0>;
|
||||
template class NPairKokkos<LMPHostType,1,0,1>;
|
||||
#endif
|
||||
}
|
||||
|
||||
@ -13,42 +13,52 @@
|
||||
|
||||
#ifdef NPAIR_CLASS
|
||||
|
||||
typedef NPairKokkos<LMPHostType,0,0> NPairKokkosFullBinHost;
|
||||
typedef NPairKokkos<LMPHostType,0,0,0> NPairKokkosFullBinHost;
|
||||
NPairStyle(full/bin/kk/host,
|
||||
NPairKokkosFullBinHost,
|
||||
NP_FULL | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPDeviceType,0,0> NPairKokkosFullBinDevice;
|
||||
typedef NPairKokkos<LMPDeviceType,0,0,0> NPairKokkosFullBinDevice;
|
||||
NPairStyle(full/bin/kk/device,
|
||||
NPairKokkosFullBinDevice,
|
||||
NP_FULL | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPHostType,0,1> NPairKokkosFullBinGhostHost;
|
||||
typedef NPairKokkos<LMPHostType,0,1,0> NPairKokkosFullBinGhostHost;
|
||||
NPairStyle(full/bin/ghost/kk/host,
|
||||
NPairKokkosFullBinGhostHost,
|
||||
NP_FULL | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPDeviceType,0,1> NPairKokkosFullBinGhostDevice;
|
||||
typedef NPairKokkos<LMPDeviceType,0,1,0> NPairKokkosFullBinGhostDevice;
|
||||
NPairStyle(full/bin/ghost/kk/device,
|
||||
NPairKokkosFullBinGhostDevice,
|
||||
NP_FULL | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPHostType,1,0> NPairKokkosHalfBinHost;
|
||||
typedef NPairKokkos<LMPHostType,1,0,0> NPairKokkosHalfBinHost;
|
||||
NPairStyle(half/bin/kk/host,
|
||||
NPairKokkosHalfBinHost,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_ORTHO)
|
||||
|
||||
typedef NPairKokkos<LMPDeviceType,1,0> NPairKokkosHalfBinDevice;
|
||||
typedef NPairKokkos<LMPDeviceType,1,0,0> NPairKokkosHalfBinDevice;
|
||||
NPairStyle(half/bin/kk/device,
|
||||
NPairKokkosHalfBinDevice,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO | NP_TRI)
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_ORTHO)
|
||||
|
||||
typedef NPairKokkos<LMPHostType,1,1> NPairKokkosHalfBinGhostHost;
|
||||
typedef NPairKokkos<LMPHostType,1,0,1> NPairKokkosHalfBinHostTri;
|
||||
NPairStyle(half/bin/kk/host,
|
||||
NPairKokkosHalfBinHostTri,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPDeviceType,1,0,1> NPairKokkosHalfBinDeviceTri;
|
||||
NPairStyle(half/bin/kk/device,
|
||||
NPairKokkosHalfBinDeviceTri,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPHostType,1,1,0> NPairKokkosHalfBinGhostHost;
|
||||
NPairStyle(half/bin/ghost/kk/host,
|
||||
NPairKokkosHalfBinGhostHost,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_HOST | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
|
||||
|
||||
typedef NPairKokkos<LMPDeviceType,1,1> NPairKokkosHalfBinGhostDevice;
|
||||
typedef NPairKokkos<LMPDeviceType,1,1,0> NPairKokkosHalfBinGhostDevice;
|
||||
NPairStyle(half/bin/ghost/kk/device,
|
||||
NPairKokkosHalfBinGhostDevice,
|
||||
NP_HALF | NP_BIN | NP_KOKKOS_DEVICE | NP_NEWTON | NP_NEWTOFF | NP_GHOST | NP_ORTHO | NP_TRI)
|
||||
@ -63,7 +73,7 @@ NPairStyle(half/bin/ghost/kk/device,
|
||||
|
||||
namespace LAMMPS_NS {
|
||||
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST>
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST, int TRI>
|
||||
class NPairKokkos : public NPair {
|
||||
public:
|
||||
NPairKokkos(class LAMMPS *);
|
||||
@ -147,7 +157,8 @@ class NeighborKokkosExecute
|
||||
// data from Atom class
|
||||
|
||||
const typename AT::t_x_array_randomread x;
|
||||
const typename AT::t_int_1d_const type,mask,molecule;
|
||||
const typename AT::t_int_1d_const type,mask;
|
||||
const typename AT::t_tagint_1d_const molecule;
|
||||
const typename AT::t_tagint_1d_const tag;
|
||||
const typename AT::t_tagint_2d_const special;
|
||||
const typename AT::t_int_2d_const nspecial;
|
||||
@ -184,7 +195,7 @@ class NeighborKokkosExecute
|
||||
const typename AT::t_x_array_randomread &_x,
|
||||
const typename AT::t_int_1d_const &_type,
|
||||
const typename AT::t_int_1d_const &_mask,
|
||||
const typename AT::t_int_1d_const &_molecule,
|
||||
const typename AT::t_tagint_1d_const &_molecule,
|
||||
const typename AT::t_tagint_1d_const &_tag,
|
||||
const typename AT::t_tagint_2d_const &_special,
|
||||
const typename AT::t_int_2d_const &_nspecial,
|
||||
@ -252,7 +263,7 @@ class NeighborKokkosExecute
|
||||
|
||||
~NeighborKokkosExecute() {neigh_list.clean_copy();};
|
||||
|
||||
template<int HalfNeigh, int Newton>
|
||||
template<int HalfNeigh, int Newton, int Tri>
|
||||
KOKKOS_FUNCTION
|
||||
void build_Item(const int &i) const;
|
||||
|
||||
@ -261,7 +272,7 @@ class NeighborKokkosExecute
|
||||
void build_Item_Ghost(const int &i) const;
|
||||
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
template<int HalfNeigh, int Newton>
|
||||
template<int HalfNeigh, int Newton, int Tri>
|
||||
__device__ inline
|
||||
void build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const;
|
||||
#endif
|
||||
@ -353,7 +364,7 @@ class NeighborKokkosExecute
|
||||
|
||||
};
|
||||
|
||||
template<class DeviceType,int HALF_NEIGH,int GHOST_NEWTON>
|
||||
template<class DeviceType, int HALF_NEIGH, int GHOST_NEWTON, int TRI>
|
||||
struct NPairKokkosBuildFunctor {
|
||||
typedef DeviceType device_type;
|
||||
|
||||
@ -366,20 +377,20 @@ struct NPairKokkosBuildFunctor {
|
||||
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
void operator() (const int & i) const {
|
||||
c.template build_Item<HALF_NEIGH,GHOST_NEWTON>(i);
|
||||
c.template build_Item<HALF_NEIGH,GHOST_NEWTON,TRI>(i);
|
||||
}
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
__device__ inline
|
||||
|
||||
void operator() (typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const {
|
||||
c.template build_ItemCuda<HALF_NEIGH,GHOST_NEWTON>(dev);
|
||||
c.template build_ItemCuda<HALF_NEIGH,GHOST_NEWTON,TRI>(dev);
|
||||
}
|
||||
size_t shmem_size(const int team_size) const { (void) team_size; return sharedsize; }
|
||||
#endif
|
||||
};
|
||||
|
||||
template<int HALF_NEIGH,int GHOST_NEWTON>
|
||||
struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {
|
||||
template<int HALF_NEIGH, int GHOST_NEWTON, int TRI>
|
||||
struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON,TRI> {
|
||||
typedef LMPHostType device_type;
|
||||
|
||||
const NeighborKokkosExecute<LMPHostType> c;
|
||||
@ -391,7 +402,7 @@ struct NPairKokkosBuildFunctor<LMPHostType,HALF_NEIGH,GHOST_NEWTON> {
|
||||
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
void operator() (const int & i) const {
|
||||
c.template build_Item<HALF_NEIGH,GHOST_NEWTON>(i);
|
||||
c.template build_Item<HALF_NEIGH,GHOST_NEWTON,TRI>(i);
|
||||
}
|
||||
|
||||
void operator() (typename Kokkos::TeamPolicy<LMPHostType>::member_type dev) const {}
|
||||
|
||||
@ -154,7 +154,7 @@ void PairBuckCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
|
||||
@ -34,6 +34,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairBuckCoulCutKokkos(class LAMMPS *);
|
||||
~PairBuckCoulCutKokkos();
|
||||
|
||||
@ -83,25 +84,25 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
@ -123,6 +124,7 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALFTHREAD,void>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairBuckCoulCutKokkos,void>(PairBuckCoulCutKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairBuckCoulCutKokkos>(PairBuckCoulCutKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -114,6 +114,19 @@ void PairBuckCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
|
||||
@ -34,6 +34,7 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairBuckCoulLongKokkos(class LAMMPS *);
|
||||
~PairBuckCoulLongKokkos();
|
||||
|
||||
@ -84,27 +85,27 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -34,6 +34,7 @@ class PairBuckKokkos : public PairBuck {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairBuckKokkos(class LAMMPS *);
|
||||
~PairBuckKokkos();
|
||||
|
||||
@ -72,22 +73,22 @@ class PairBuckKokkos : public PairBuck {
|
||||
typename Kokkos::DualView<params_buck**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_buck m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_tagint_1d tag;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
|
||||
@ -83,6 +83,19 @@ void PairCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -124,6 +137,16 @@ void PairCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
}
|
||||
|
||||
|
||||
@ -34,6 +34,7 @@ class PairCoulCutKokkos : public PairCoulCut {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairCoulCutKokkos(class LAMMPS *);
|
||||
~PairCoulCutKokkos();
|
||||
|
||||
@ -87,22 +88,25 @@ class PairCoulCutKokkos : public PairCoulCut {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
|
||||
@ -90,6 +90,19 @@ void PairCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -136,7 +149,17 @@ void PairCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -34,6 +34,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairCoulDebyeKokkos(class LAMMPS *);
|
||||
~PairCoulDebyeKokkos();
|
||||
|
||||
@ -87,22 +88,25 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
@ -124,6 +128,7 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALFTHREAD,void>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairCoulDebyeKokkos,void>(PairCoulDebyeKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairCoulDebyeKokkos>(PairCoulDebyeKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -88,12 +88,12 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -183,8 +183,6 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -195,6 +193,8 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
@ -37,6 +37,7 @@ class PairCoulDSFKokkos : public PairCoulDSF {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
typedef EV_FLOAT value_type;
|
||||
PairCoulDSFKokkos(class LAMMPS *);
|
||||
~PairCoulDSFKokkos();
|
||||
@ -61,16 +62,16 @@ class PairCoulDSFKokkos : public PairCoulDSF {
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
int sbmask(const int& j) const;
|
||||
|
||||
protected:
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
protected:
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
|
||||
int neighflag,newton_pair;
|
||||
@ -79,9 +80,9 @@ class PairCoulDSFKokkos : public PairCoulDSF {
|
||||
double special_coul[4];
|
||||
double qqrd2e;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
friend void pair_virial_fdotr_compute<PairCoulDSFKokkos>(PairCoulDSFKokkos*);
|
||||
|
||||
@ -107,6 +107,19 @@ void PairCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_coulsq.template sync<DeviceType>();
|
||||
@ -158,6 +171,16 @@ void PairCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairCoulLongKokkos : public PairCoulLong {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairCoulLongKokkos(class LAMMPS *);
|
||||
~PairCoulLongKokkos();
|
||||
|
||||
@ -45,7 +46,9 @@ class PairCoulLongKokkos : public PairCoulLong {
|
||||
double init_one(int, int);
|
||||
|
||||
struct params_coul{
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
params_coul(){cut_coulsq=0;};
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
params_coul(int i){cut_coulsq=0;};
|
||||
F_FLOAT cut_coulsq;
|
||||
};
|
||||
@ -86,27 +89,27 @@ class PairCoulLongKokkos : public PairCoulLong {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -83,12 +83,12 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -184,8 +184,6 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -196,6 +194,8 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
@ -37,6 +37,7 @@ class PairCoulWolfKokkos : public PairCoulWolf {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
typedef EV_FLOAT value_type;
|
||||
PairCoulWolfKokkos(class LAMMPS *);
|
||||
~PairCoulWolfKokkos();
|
||||
@ -63,14 +64,14 @@ class PairCoulWolfKokkos : public PairCoulWolf {
|
||||
|
||||
protected:
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
|
||||
int neighflag,newton_pair;
|
||||
@ -81,9 +82,9 @@ class PairCoulWolfKokkos : public PairCoulWolf {
|
||||
double special_coul[4];
|
||||
double qqrd2e;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
friend void pair_virial_fdotr_compute<PairCoulWolfKokkos>(PairCoulWolfKokkos*);
|
||||
|
||||
@ -82,12 +82,12 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -246,8 +246,6 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -258,6 +256,8 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
@ -125,8 +125,8 @@ class PairEAMAlloyKokkos : public PairEAM {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
DAT::tdual_ffloat_1d k_rho;
|
||||
DAT::tdual_ffloat_1d k_fp;
|
||||
@ -154,9 +154,9 @@ class PairEAMAlloyKokkos : public PairEAM {
|
||||
void interpolate(int, double, double *, t_host_ffloat_2d_n7, int);
|
||||
void read_file(char *);
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int iswap;
|
||||
|
||||
@ -82,12 +82,12 @@ void PairEAMFSKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
|
||||
@ -125,8 +125,8 @@ class PairEAMFSKokkos : public PairEAM {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
DAT::tdual_ffloat_1d k_rho;
|
||||
DAT::tdual_ffloat_1d k_fp;
|
||||
@ -154,9 +154,9 @@ class PairEAMFSKokkos : public PairEAM {
|
||||
void interpolate(int, double, double *, t_host_ffloat_2d_n7, int);
|
||||
void read_file(char *);
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int iswap;
|
||||
|
||||
@ -77,12 +77,12 @@ void PairEAMKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
|
||||
@ -122,8 +122,8 @@ class PairEAMKokkos : public PairEAM {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
DAT::tdual_ffloat_1d k_rho;
|
||||
DAT::tdual_ffloat_1d k_fp;
|
||||
@ -149,9 +149,9 @@ class PairEAMKokkos : public PairEAM {
|
||||
virtual void file2array();
|
||||
void array2spline();
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int iswap;
|
||||
|
||||
@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::compute(int eflag_in, int
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::compute(int eflag_in, int
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCharmmCoulCharmmImplicitKokkos(class LAMMPS *);
|
||||
~PairLJCharmmCoulCharmmImplicitKokkos();
|
||||
|
||||
@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -115,6 +115,19 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_i
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -167,6 +180,16 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_i
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCharmmCoulCharmmKokkos(class LAMMPS *);
|
||||
~PairLJCharmmCoulCharmmKokkos();
|
||||
|
||||
@ -77,27 +78,27 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -180,8 +180,6 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -192,6 +190,8 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCharmmCoulLongKokkos(class LAMMPS *);
|
||||
~PairLJCharmmCoulLongKokkos();
|
||||
|
||||
@ -75,27 +76,27 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -92,6 +92,19 @@ void PairLJClass2CoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -138,7 +151,17 @@ void PairLJClass2CoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJClass2CoulCutKokkos(class LAMMPS *);
|
||||
~PairLJClass2CoulCutKokkos();
|
||||
|
||||
@ -76,22 +77,25 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
int neighflag;
|
||||
int nlocal,nall,eflag,vflag;
|
||||
@ -112,6 +116,7 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALFTHREAD,void>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairLJClass2CoulCutKokkos,void>(PairLJClass2CoulCutKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairLJClass2CoulCutKokkos>(PairLJClass2CoulCutKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -100,6 +100,19 @@ void PairLJClass2CoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -152,6 +165,16 @@ void PairLJClass2CoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJClass2CoulLongKokkos(class LAMMPS *);
|
||||
~PairLJClass2CoulLongKokkos();
|
||||
|
||||
@ -76,24 +77,27 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -92,6 +92,19 @@ void PairLJClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_params.template sync<DeviceType>();
|
||||
@ -125,7 +138,18 @@ void PairLJClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJClass2Kokkos : public PairLJClass2 {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJClass2Kokkos(class LAMMPS *);
|
||||
~PairLJClass2Kokkos();
|
||||
|
||||
@ -79,19 +80,22 @@ class PairLJClass2Kokkos : public PairLJClass2 {
|
||||
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_tagint_1d tag;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
int neighflag;
|
||||
int nlocal,nall,eflag,vflag;
|
||||
|
||||
@ -92,6 +92,19 @@ void PairLJCutCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -136,7 +149,18 @@ void PairLJCutCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCutCoulCutKokkos(class LAMMPS *);
|
||||
~PairLJCutCoulCutKokkos();
|
||||
|
||||
@ -75,22 +76,25 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
@ -112,6 +116,7 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALFTHREAD,void>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairLJCutCoulCutKokkos,void>(PairLJCutCoulCutKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairLJCutCoulCutKokkos>(PairLJCutCoulCutKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -96,6 +96,19 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -142,7 +155,17 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCutCoulDebyeKokkos(class LAMMPS *);
|
||||
~PairLJCutCoulDebyeKokkos();
|
||||
|
||||
@ -75,22 +76,25 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
@ -112,6 +116,7 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALFTHREAD,void>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairLJCutCoulDebyeKokkos,void>(PairLJCutCoulDebyeKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairLJCutCoulDebyeKokkos>(PairLJCutCoulDebyeKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -104,6 +104,19 @@ void PairLJCutCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -160,7 +173,17 @@ void PairLJCutCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) virial_fdotr_compute();
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCutCoulDSFKokkos(class LAMMPS *);
|
||||
~PairLJCutCoulDSFKokkos();
|
||||
|
||||
@ -74,22 +75,25 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
@ -111,6 +115,7 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
|
||||
friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALFTHREAD,void>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairLJCutCoulDSFKokkos,void>(PairLJCutCoulDSFKokkos*,
|
||||
NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairLJCutCoulDSFKokkos>(PairLJCutCoulDSFKokkos*);
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -168,8 +168,6 @@ void PairLJCutCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -179,6 +177,9 @@ void PairLJCutCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCutCoulLongKokkos(class LAMMPS *);
|
||||
~PairLJCutCoulLongKokkos();
|
||||
|
||||
@ -76,27 +77,27 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -138,8 +138,6 @@ void PairLJCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -149,6 +147,9 @@ void PairLJCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJCutKokkos : public PairLJCut {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJCutKokkos(class LAMMPS *);
|
||||
~PairLJCutKokkos();
|
||||
|
||||
@ -73,22 +74,22 @@ class PairLJCutKokkos : public PairLJCut {
|
||||
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
typename AT::t_tagint_1d tag;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
|
||||
@ -91,6 +91,19 @@ void PairLJExpandKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_params.template sync<DeviceType>();
|
||||
@ -126,6 +139,16 @@ void PairLJExpandKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJExpandKokkos : public PairLJExpand {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJExpandKokkos(class LAMMPS *);
|
||||
~PairLJExpandKokkos();
|
||||
|
||||
@ -79,19 +80,22 @@ class PairLJExpandKokkos : public PairLJExpand {
|
||||
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_tagint_1d tag;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
|
||||
@ -106,6 +106,19 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_ljsq.template sync<DeviceType>();
|
||||
@ -158,6 +171,16 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=1};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJGromacsCoulGromacsKokkos(class LAMMPS *);
|
||||
~PairLJGromacsCoulGromacsKokkos();
|
||||
|
||||
@ -84,27 +85,27 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_ljsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_coulsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_ljsq;
|
||||
typename AT::t_ffloat_2d d_cut_ljsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_coulsq;
|
||||
typename AT::t_ffloat_2d d_cut_coulsq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -103,6 +103,19 @@ void PairLJGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_cut_inner.template sync<DeviceType>();
|
||||
@ -145,6 +158,16 @@ void PairLJGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJGromacsKokkos : public PairLJGromacs {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJGromacsKokkos(class LAMMPS *);
|
||||
~PairLJGromacsKokkos();
|
||||
|
||||
@ -87,27 +88,27 @@ class PairLJGromacsKokkos : public PairLJGromacs {
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_inner[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cut_inner_sq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_float_1d_randomread q;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_float_1d_randomread q;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_inner;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_inner;
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cut_inner_sq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cut_inner_sq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cut_inner;
|
||||
typename AT::t_ffloat_2d d_cut_inner;
|
||||
typename AT::tdual_ffloat_2d k_cut_inner_sq;
|
||||
typename AT::t_ffloat_2d d_cut_inner_sq;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread
|
||||
typename AT::t_ffloat_1d_randomread
|
||||
d_rtable, d_drtable, d_ftable, d_dftable,
|
||||
d_ctable, d_dctable, d_etable, d_detable;
|
||||
|
||||
|
||||
@ -91,6 +91,19 @@ void PairLJSDKKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_params.template sync<DeviceType>();
|
||||
@ -124,7 +137,18 @@ void PairLJSDKKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
|
||||
@ -34,6 +34,7 @@ class PairLJSDKKokkos : public PairLJSDK {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
PairLJSDKKokkos(class LAMMPS *);
|
||||
~PairLJSDKKokkos();
|
||||
|
||||
@ -74,19 +75,22 @@ class PairLJSDKKokkos : public PairLJSDK {
|
||||
typename Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_lj m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; // hardwired to space for 12 atom types
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
typename AT::t_tagint_1d tag;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::tdual_ffloat_2d k_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
|
||||
310
src/KOKKOS/pair_morse_kokkos.cpp
Normal file
310
src/KOKKOS/pair_morse_kokkos.cpp
Normal file
@ -0,0 +1,310 @@
|
||||
/* ----------------------------------------------------------------------
|
||||
LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
|
||||
http://lammps.sandia.gov, Sandia National Laboratories
|
||||
Steve Plimpton, sjplimp@sandia.gov
|
||||
|
||||
Copyright (2003) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
certain rights in this software. This software is distributed under
|
||||
the GNU General Public License.
|
||||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing authors: Stefan Paquay (Eindhoven University of Technology)
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
#include <math.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include "pair_morse_kokkos.h"
|
||||
#include "kokkos.h"
|
||||
#include "atom_kokkos.h"
|
||||
#include "comm.h"
|
||||
#include "force.h"
|
||||
#include "neighbor.h"
|
||||
#include "neigh_list.h"
|
||||
#include "neigh_request.h"
|
||||
#include "update.h"
|
||||
#include "integrate.h"
|
||||
#include "respa.h"
|
||||
#include "math_const.h"
|
||||
#include "memory.h"
|
||||
#include "error.h"
|
||||
#include "atom_masks.h"
|
||||
|
||||
using namespace LAMMPS_NS;
|
||||
using namespace MathConst;
|
||||
|
||||
#define KOKKOS_CUDA_MAX_THREADS 256
|
||||
#define KOKKOS_CUDA_MIN_BLOCKS 8
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
PairMorseKokkos<DeviceType>::PairMorseKokkos(LAMMPS *lmp) : PairMorse(lmp)
|
||||
{
|
||||
respa_enable = 0;
|
||||
|
||||
atomKK = (AtomKokkos *) atom;
|
||||
execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
|
||||
datamask_read = X_MASK | F_MASK | TYPE_MASK | ENERGY_MASK | VIRIAL_MASK;
|
||||
datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK;
|
||||
cutsq = NULL;
|
||||
}
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
PairMorseKokkos<DeviceType>::~PairMorseKokkos()
|
||||
{
|
||||
if (allocated) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
k_cutsq = DAT::tdual_ffloat_2d();
|
||||
memory->sfree(cutsq);
|
||||
eatom = NULL;
|
||||
vatom = NULL;
|
||||
cutsq = NULL;
|
||||
}
|
||||
}
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
void PairMorseKokkos<DeviceType>::cleanup_copy() {
|
||||
// WHY needed: this prevents parent copy from deallocating any arrays
|
||||
allocated = 0;
|
||||
cutsq = NULL;
|
||||
eatom = NULL;
|
||||
vatom = NULL;
|
||||
}
|
||||
|
||||
/* ---------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
void PairMorseKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
{
|
||||
eflag = eflag_in;
|
||||
vflag = vflag_in;
|
||||
|
||||
|
||||
if (neighflag == FULL) no_virial_fdotr_compute = 1;
|
||||
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
k_cutsq.template sync<DeviceType>();
|
||||
k_params.template sync<DeviceType>();
|
||||
if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
|
||||
else atomKK->modified(execution_space,F_MASK);
|
||||
|
||||
x = atomKK->k_x.view<DeviceType>();
|
||||
c_x = atomKK->k_x.view<DeviceType>();
|
||||
f = atomKK->k_f.view<DeviceType>();
|
||||
type = atomKK->k_type.view<DeviceType>();
|
||||
tag = atomKK->k_tag.view<DeviceType>();
|
||||
nlocal = atom->nlocal;
|
||||
nall = atom->nlocal + atom->nghost;
|
||||
newton_pair = force->newton_pair;
|
||||
special_lj[0] = force->special_lj[0];
|
||||
special_lj[1] = force->special_lj[1];
|
||||
special_lj[2] = force->special_lj[2];
|
||||
special_lj[3] = force->special_lj[3];
|
||||
|
||||
// loop over neighbors of my atoms
|
||||
|
||||
EV_FLOAT ev = pair_compute<PairMorseKokkos<DeviceType>,void >(this,(NeighListKokkos<DeviceType>*)list);
|
||||
|
||||
if (eflag_global) eng_vdwl += ev.evdwl;
|
||||
if (vflag_global) {
|
||||
virial[0] += ev.v[0];
|
||||
virial[1] += ev.v[1];
|
||||
virial[2] += ev.v[2];
|
||||
virial[3] += ev.v[3];
|
||||
virial[4] += ev.v[4];
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
template<bool STACKPARAMS, class Specialisation>
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
F_FLOAT PairMorseKokkos<DeviceType>::
|
||||
compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
|
||||
(void) i;
|
||||
(void) j;
|
||||
const F_FLOAT rr = sqrt(rsq);
|
||||
const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0;
|
||||
const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0;
|
||||
const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha;
|
||||
const F_FLOAT dr = rr - r0;
|
||||
|
||||
// U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ]
|
||||
// f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r)
|
||||
// = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r)
|
||||
const F_FLOAT dexp = exp( -aa*dr );
|
||||
const F_FLOAT forcelj = 2*aa*d0*dexp*(dexp-1.0);
|
||||
|
||||
return forcelj / rr;
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
template<bool STACKPARAMS, class Specialisation>
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
F_FLOAT PairMorseKokkos<DeviceType>::
|
||||
compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
|
||||
(void) i;
|
||||
(void) j;
|
||||
const F_FLOAT rr = sqrt(rsq);
|
||||
const F_FLOAT r0 = STACKPARAMS ? m_params[itype][jtype].r0 : params(itype,jtype).r0;
|
||||
const F_FLOAT d0 = STACKPARAMS ? m_params[itype][jtype].d0 : params(itype,jtype).d0;
|
||||
const F_FLOAT aa = STACKPARAMS ? m_params[itype][jtype].alpha : params(itype,jtype).alpha;
|
||||
const F_FLOAT dr = rr - r0;
|
||||
|
||||
// U = d0 * [ exp( -2*a*(x-r0)) - 2*exp(-a*(x-r0)) ]
|
||||
// f = -2*a*d0*[ -exp( -2*a*(x-r0) ) + exp( -a*(x-r0) ) ] * grad(r)
|
||||
// = +2*a*d0*[ exp( -2*a*(x-r0) ) - exp( -a*(x-r0) ) ] * grad(r)
|
||||
const F_FLOAT dexp = exp( -aa*dr );
|
||||
|
||||
return d0 * dexp * ( dexp - 2.0 );
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
allocate all arrays
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
void PairMorseKokkos<DeviceType>::allocate()
|
||||
{
|
||||
PairMorse::allocate();
|
||||
|
||||
int n = atom->ntypes;
|
||||
memory->destroy(cutsq);
|
||||
memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq");
|
||||
d_cutsq = k_cutsq.template view<DeviceType>();
|
||||
k_params = Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType>("PairMorse::params",n+1,n+1);
|
||||
params = k_params.d_view;
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
global settings
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
void PairMorseKokkos<DeviceType>::settings(int narg, char **arg)
|
||||
{
|
||||
if (narg > 2) error->all(FLERR,"Illegal pair_style command");
|
||||
|
||||
PairMorse::settings(1,arg);
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
init specific to this pair style
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
template<class DeviceType>
|
||||
void PairMorseKokkos<DeviceType>::init_style()
|
||||
{
|
||||
PairMorse::init_style();
|
||||
|
||||
// error if rRESPA with inner levels
|
||||
|
||||
if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) {
|
||||
int respa = 0;
|
||||
if (((Respa *) update->integrate)->level_inner >= 0) respa = 1;
|
||||
if (((Respa *) update->integrate)->level_middle >= 0) respa = 2;
|
||||
if (respa)
|
||||
error->all(FLERR,"Cannot use Kokkos pair style with rRESPA inner/middle");
|
||||
}
|
||||
|
||||
// irequest = neigh request made by parent class
|
||||
|
||||
neighflag = lmp->kokkos->neighflag;
|
||||
int irequest = neighbor->nrequest - 1;
|
||||
|
||||
neighbor->requests[irequest]->
|
||||
kokkos_host = Kokkos::Impl::is_same<DeviceType,LMPHostType>::value &&
|
||||
!Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
|
||||
neighbor->requests[irequest]->
|
||||
kokkos_device = Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
|
||||
|
||||
if (neighflag == FULL) {
|
||||
neighbor->requests[irequest]->full = 1;
|
||||
neighbor->requests[irequest]->half = 0;
|
||||
} else if (neighflag == HALF || neighflag == HALFTHREAD) {
|
||||
neighbor->requests[irequest]->full = 0;
|
||||
neighbor->requests[irequest]->half = 1;
|
||||
} else if (neighflag == N2) {
|
||||
neighbor->requests[irequest]->full = 0;
|
||||
neighbor->requests[irequest]->half = 0;
|
||||
} else {
|
||||
error->all(FLERR,"Cannot use chosen neighbor list style with morse/kk");
|
||||
}
|
||||
}
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
init for one type pair i,j and corresponding j,i
|
||||
------------------------------------------------------------------------- */
|
||||
// Rewrite this.
|
||||
template<class DeviceType>
|
||||
double PairMorseKokkos<DeviceType>::init_one(int i, int j)
|
||||
{
|
||||
double cutone = PairMorse::init_one(i,j);
|
||||
|
||||
k_params.h_view(i,j).d0 = d0[i][j];
|
||||
k_params.h_view(i,j).alpha = alpha[i][j];
|
||||
k_params.h_view(i,j).r0 = r0[i][j];
|
||||
k_params.h_view(i,j).offset = offset[i][j];
|
||||
k_params.h_view(i,j).cutsq = cutone*cutone;
|
||||
k_params.h_view(j,i) = k_params.h_view(i,j);
|
||||
|
||||
if(i<MAX_TYPES_STACKPARAMS+1 && j<MAX_TYPES_STACKPARAMS+1) {
|
||||
m_params[i][j] = m_params[j][i] = k_params.h_view(i,j);
|
||||
m_cutsq[j][i] = m_cutsq[i][j] = cutone*cutone;
|
||||
}
|
||||
|
||||
k_cutsq.h_view(i,j) = k_cutsq.h_view(j,i) = cutone*cutone;
|
||||
k_cutsq.template modify<LMPHostType>();
|
||||
k_params.template modify<LMPHostType>();
|
||||
|
||||
return cutone;
|
||||
}
|
||||
|
||||
|
||||
|
||||
namespace LAMMPS_NS {
|
||||
template class PairMorseKokkos<LMPDeviceType>;
|
||||
#ifdef KOKKOS_HAVE_CUDA
|
||||
template class PairMorseKokkos<LMPHostType>;
|
||||
#endif
|
||||
}
|
||||
|
||||
135
src/KOKKOS/pair_morse_kokkos.h
Normal file
135
src/KOKKOS/pair_morse_kokkos.h
Normal file
@ -0,0 +1,135 @@
|
||||
/* -*- c++ -*- ----------------------------------------------------------
|
||||
LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
|
||||
http://lammps.sandia.gov, Sandia National Laboratories
|
||||
Steve Plimpton, sjplimp@sandia.gov
|
||||
|
||||
Copyright (2003) Sandia Corporation. Under the terms of Contract
|
||||
DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
|
||||
certain rights in this software. This software is distributed under
|
||||
the GNU General Public License.
|
||||
|
||||
See the README file in the top-level LAMMPS directory.
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
#ifdef PAIR_CLASS
|
||||
|
||||
PairStyle(morse/kk,PairMorseKokkos<LMPDeviceType>)
|
||||
PairStyle(morse/kk/device,PairMorseKokkos<LMPDeviceType>)
|
||||
PairStyle(morse/kk/host,PairMorseKokkos<LMPHostType>)
|
||||
|
||||
#else
|
||||
|
||||
#ifndef LMP_PAIR_MORSE_KOKKOS_H
|
||||
#define LMP_PAIR_MORSE_KOKKOS_H
|
||||
|
||||
#include "pair_kokkos.h"
|
||||
#include "pair_morse.h"
|
||||
#include "neigh_list_kokkos.h"
|
||||
|
||||
namespace LAMMPS_NS {
|
||||
|
||||
template<class DeviceType>
|
||||
class PairMorseKokkos : public PairMorse {
|
||||
public:
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
PairMorseKokkos(class LAMMPS *);
|
||||
virtual ~PairMorseKokkos();
|
||||
|
||||
void compute(int, int);
|
||||
|
||||
void settings(int, char **);
|
||||
void init_style();
|
||||
double init_one(int, int);
|
||||
|
||||
struct params_morse{
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
params_morse(){cutsq=0,d0=0;alpha=0;r0=0;offset=0;}
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
params_morse(int i){cutsq=0,d0=0;alpha=0;r0=0;offset=0;}
|
||||
F_FLOAT cutsq,d0,alpha,r0,offset;
|
||||
};
|
||||
|
||||
protected:
|
||||
void cleanup_copy();
|
||||
|
||||
template<bool STACKPARAMS, class Specialisation>
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const;
|
||||
|
||||
template<bool STACKPARAMS, class Specialisation>
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const;
|
||||
|
||||
template<bool STACKPARAMS, class Specialisation>
|
||||
KOKKOS_INLINE_FUNCTION
|
||||
F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const {
|
||||
return 0;
|
||||
}
|
||||
|
||||
|
||||
Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType> k_params;
|
||||
typename Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType>::t_dev_const_um params;
|
||||
params_morse m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
|
||||
int newton_pair;
|
||||
double special_lj[4];
|
||||
|
||||
typename ArrayTypes<DeviceType>::tdual_ffloat_2d k_cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
|
||||
|
||||
int neighflag;
|
||||
int nlocal,nall,eflag,vflag;
|
||||
|
||||
void allocate();
|
||||
friend class PairComputeFunctor<PairMorseKokkos,FULL,true>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,HALF,true>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,HALFTHREAD,true>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,N2,true>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,FULL,false>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,HALF,false>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,HALFTHREAD,false>;
|
||||
friend class PairComputeFunctor<PairMorseKokkos,N2,false>;
|
||||
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,FULL,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALF,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALFTHREAD,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,N2,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend EV_FLOAT pair_compute<PairMorseKokkos,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
|
||||
friend void pair_virial_fdotr_compute<PairMorseKokkos>(PairMorseKokkos*);
|
||||
};
|
||||
|
||||
}
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
/* ERROR/WARNING messages:
|
||||
|
||||
E: Illegal ... command
|
||||
|
||||
Self-explanatory. Check the input script syntax and compare to the
|
||||
documentation for the command. You can use -echo screen as a
|
||||
command-line option when running LAMMPS to see the offending line.
|
||||
|
||||
E: Cannot use Kokkos pair style with rRESPA inner/middle
|
||||
|
||||
Self-explanatory.
|
||||
|
||||
E: Cannot use chosen neighbor list style with morse/kk
|
||||
|
||||
That style is not supported by Kokkos.
|
||||
|
||||
*/
|
||||
@ -1046,7 +1046,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeLJCoulomb<NEIGHFLAG,
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const F_FLOAT qi = q(i);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int jnum = d_numneigh[i];
|
||||
|
||||
F_FLOAT fxtmp, fytmp, fztmp;
|
||||
@ -1056,7 +1056,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeLJCoulomb<NEIGHFLAG,
|
||||
int j = d_neighbors(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtype = type(j);
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
const F_FLOAT qj = q(j);
|
||||
|
||||
if (NEIGHFLAG != FULL) {
|
||||
@ -1201,7 +1201,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTabulatedLJCoulomb<N
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const F_FLOAT qi = q(i);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int jnum = d_numneigh[i];
|
||||
|
||||
F_FLOAT fxtmp, fytmp, fztmp;
|
||||
@ -1211,7 +1211,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTabulatedLJCoulomb<N
|
||||
int j = d_neighbors(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtype = type(j);
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
const F_FLOAT qj = q(j);
|
||||
|
||||
if (NEIGHFLAG != FULL) {
|
||||
@ -1580,7 +1580,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf<NEIGHFLAG>,
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int jnum = d_numneigh[i];
|
||||
|
||||
F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3], dBOp_i[3], dln_BOp_pi_i[3], dln_BOp_pi2_i[3];
|
||||
@ -1605,7 +1605,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf<NEIGHFLAG>,
|
||||
for (int jj = 0; jj < jnum; jj++) {
|
||||
int j = d_neighbors(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
d_bo_first[j] = j*maxbo;
|
||||
d_hb_first[j] = j*maxhb;
|
||||
@ -1802,7 +1802,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf_LessAtomics<
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int jnum = d_numneigh[i];
|
||||
|
||||
F_FLOAT C12, C34, C56, BO_s, BO_pi, BO_pi2, BO, delij[3];
|
||||
@ -1826,7 +1826,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxBuildListsHalf_LessAtomics<
|
||||
for (int jj = 0; jj < jnum; jj++) {
|
||||
int j = d_neighbors(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
d_bo_first[j] = j*maxbo;
|
||||
d_hb_first[j] = j*maxhb;
|
||||
@ -2752,7 +2752,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTorsion<NEIGHFLAG,EV
|
||||
|
||||
const int i = d_ilist[ii];
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const X_FLOAT xtmp = x(i,0);
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
@ -2768,7 +2768,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeTorsion<NEIGHFLAG,EV
|
||||
for (int jj = j_start; jj < j_end; jj++) {
|
||||
int j = d_bo_list[jj];
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
const int jtype = type(j);
|
||||
const int j_index = jj - j_start;
|
||||
|
||||
@ -3113,7 +3113,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
|
||||
const X_FLOAT xtmp = x(i,0);
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
|
||||
const int j_start = d_bo_first[i];
|
||||
const int j_end = j_start + d_bo_num[i];
|
||||
@ -3140,7 +3140,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
|
||||
for (int kk = k_start; kk < k_end; kk++) {
|
||||
int k = d_hb_list[kk];
|
||||
k &= NEIGHMASK;
|
||||
const int ktag = tag(k);
|
||||
const tagint ktag = tag(k);
|
||||
const int ktype = type(k);
|
||||
|
||||
delik[0] = x(k,0) - xtmp;
|
||||
@ -3153,7 +3153,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeHydrogen<NEIGHFLAG,E
|
||||
const int jj = hblist[itr];
|
||||
int j = d_bo_list[jj];
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
if (jtag == ktag) continue;
|
||||
|
||||
const int jtype = type(j);
|
||||
@ -3247,14 +3247,14 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxUpdateBond<NEIGHFLAG>, cons
|
||||
Kokkos::View<F_FLOAT**, typename DAT::t_ffloat_2d_dl::array_layout,DeviceType,Kokkos::MemoryTraits<AtomicF<NEIGHFLAG>::value> > a_Cdbopi2 = d_Cdbopi2;
|
||||
|
||||
const int i = d_ilist[ii];
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const int j_start = d_bo_first[i];
|
||||
const int j_end = j_start + d_bo_num[i];
|
||||
|
||||
for (int jj = j_start; jj < j_end; jj++) {
|
||||
int j = d_bo_list[jj];
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
const int j_index = jj - j_start;
|
||||
const F_FLOAT Cdbo_i = d_Cdbo(i,j_index);
|
||||
const F_FLOAT Cdbopi_i = d_Cdbopi(i,j_index);
|
||||
@ -3302,7 +3302,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond1<NEIGHFLAG,EVFL
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const F_FLOAT imass = paramssing(itype).mass;
|
||||
const F_FLOAT val_i = paramssing(itype).valency;
|
||||
const int j_start = d_bo_first[i];
|
||||
@ -3313,7 +3313,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond1<NEIGHFLAG,EVFL
|
||||
for (int jj = j_start; jj < j_end; jj++) {
|
||||
int j = d_bo_list[jj];
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
if (itag > jtag) {
|
||||
if ((itag+jtag) % 2 == 0) continue;
|
||||
@ -3438,7 +3438,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond2<NEIGHFLAG,EVFL
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
const F_FLOAT imass = paramssing(itype).mass;
|
||||
const F_FLOAT val_i = paramssing(itype).valency;
|
||||
const int j_start = d_bo_first[i];
|
||||
@ -3451,7 +3451,7 @@ void PairReaxCKokkos<DeviceType>::operator()(PairReaxComputeBond2<NEIGHFLAG,EVFL
|
||||
for (int jj = j_start; jj < j_end; jj++) {
|
||||
int j = d_bo_list[jj];
|
||||
j &= NEIGHMASK;
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
if (itag > jtag) {
|
||||
if ((itag+jtag) % 2 == 0) continue;
|
||||
|
||||
@ -392,7 +392,7 @@ class PairReaxCKokkos : public PairReaxC {
|
||||
typename AT::t_efloat_1d v_eatom;
|
||||
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_virial_array v_vatom;
|
||||
HAT::t_virial_array h_vatom;
|
||||
|
||||
@ -414,9 +414,9 @@ class PairReaxCKokkos : public PairReaxC {
|
||||
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
|
||||
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
|
||||
typename AT::t_int_1d d_bo_first, d_bo_num, d_bo_list, d_hb_first, d_hb_num, d_hb_list;
|
||||
|
||||
|
||||
@ -88,12 +88,12 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -171,8 +171,6 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev_all.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -183,6 +181,8 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
@ -118,23 +118,23 @@ class PairSWKokkos : public PairSW {
|
||||
void threebodyj(const Param&, const Param&, const Param&, const F_FLOAT&, const F_FLOAT&, F_FLOAT *, F_FLOAT *,
|
||||
F_FLOAT *) const;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_tagint_1d tag;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_tagint_1d tag;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
DAT::t_int_1d_randomread d_type2frho;
|
||||
DAT::t_int_2d_randomread d_type2rhor;
|
||||
DAT::t_int_2d_randomread d_type2z2r;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int neighflag,newton_pair;
|
||||
|
||||
@ -12,7 +12,7 @@
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing author: Paul Crozier (SNL)
|
||||
Contributing author: Christian Trott (SNL)
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
#include <mpi.h>
|
||||
@ -89,6 +89,19 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
|
||||
if (eflag || vflag) ev_setup(eflag,vflag);
|
||||
else evflag = vflag_fdotr = 0;
|
||||
|
||||
// reallocate per-atom arrays if necessary
|
||||
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
//k_cutsq.template sync<DeviceType>();
|
||||
//k_params.template sync<DeviceType>();
|
||||
@ -165,7 +178,18 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
|
||||
virial[5] += ev.v[5];
|
||||
}
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_atom) {
|
||||
k_vatom.template modify<DeviceType>();
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
}
|
||||
|
||||
template<class DeviceType>
|
||||
|
||||
@ -44,6 +44,7 @@ class PairTableKokkos : public PairTable {
|
||||
enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF|N2};
|
||||
enum {COUL_FLAG=0};
|
||||
typedef DeviceType device_type;
|
||||
typedef ArrayTypes<DeviceType> AT;
|
||||
|
||||
PairTableKokkos(class LAMMPS *);
|
||||
virtual ~PairTableKokkos();
|
||||
@ -62,27 +63,27 @@ class PairTableKokkos : public PairTable {
|
||||
protected:
|
||||
|
||||
/*struct TableDeviceConst {
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_int_2d_randomread tabindex;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread nshiftbits,nmask;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d_randomread innersq,invdelta,deltasq6;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
|
||||
typename AT::t_ffloat_2d_randomread cutsq;
|
||||
typename AT::t_int_2d_randomread tabindex;
|
||||
typename AT::t_int_1d_randomread nshiftbits,nmask;
|
||||
typename AT::t_ffloat_1d_randomread innersq,invdelta,deltasq6;
|
||||
typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
|
||||
};*/
|
||||
//Its faster not to use texture fetch if the number of tables is less than 32!
|
||||
struct TableDeviceConst {
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_int_2d tabindex;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d nshiftbits,nmask;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d innersq,invdelta,deltasq6;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
|
||||
typename AT::t_ffloat_2d cutsq;
|
||||
typename AT::t_int_2d tabindex;
|
||||
typename AT::t_int_1d nshiftbits,nmask;
|
||||
typename AT::t_ffloat_1d innersq,invdelta,deltasq6;
|
||||
typename AT::t_ffloat_2d_randomread rsq,drsq,e,de,f,df,e2,f2;
|
||||
};
|
||||
|
||||
struct TableDevice {
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d cutsq;
|
||||
typename ArrayTypes<DeviceType>::t_int_2d tabindex;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d nshiftbits,nmask;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_1d innersq,invdelta,deltasq6;
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2;
|
||||
typename AT::t_ffloat_2d cutsq;
|
||||
typename AT::t_int_2d tabindex;
|
||||
typename AT::t_int_1d nshiftbits,nmask;
|
||||
typename AT::t_ffloat_1d innersq,invdelta,deltasq6;
|
||||
typename AT::t_ffloat_2d rsq,drsq,e,de,f,df,e2,f2;
|
||||
};
|
||||
|
||||
struct TableHost {
|
||||
@ -99,17 +100,20 @@ class PairTableKokkos : public PairTable {
|
||||
|
||||
F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1];
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_ffloat_2d d_cutsq;
|
||||
typename AT::t_ffloat_2d d_cutsq;
|
||||
|
||||
virtual void allocate();
|
||||
void compute_table(Table *);
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_x_array_randomread x;
|
||||
typename ArrayTypes<DeviceType>::t_x_array_const c_x;
|
||||
typename ArrayTypes<DeviceType>::t_f_array f;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread type;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
typename AT::t_x_array_randomread x;
|
||||
typename AT::t_x_array_const c_x;
|
||||
typename AT::t_f_array f;
|
||||
typename AT::t_int_1d_randomread type;
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
typename AT::t_efloat_1d d_eatom;
|
||||
typename AT::t_virial_array d_vatom;
|
||||
|
||||
protected:
|
||||
int nlocal,nall,eflag,vflag,neighflag,newton_pair;
|
||||
|
||||
@ -172,12 +172,12 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -254,8 +254,6 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev_all.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -266,6 +264,8 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -314,7 +314,7 @@ void PairTersoffKokkos<DeviceType>::operator()(TagPairTersoffComputeHalf<NEIGHFL
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
|
||||
F_FLOAT fi[3], fj[3], fk[3];
|
||||
|
||||
@ -331,7 +331,7 @@ void PairTersoffKokkos<DeviceType>::operator()(TagPairTersoffComputeHalf<NEIGHFL
|
||||
int j = d_neighbors_short(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtype = type(j);
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
if (itag > jtag) {
|
||||
if ((itag+jtag) % 2 == 0) continue;
|
||||
|
||||
@ -197,16 +197,16 @@ class PairTersoffKokkos : public PairTersoff {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
|
||||
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
|
||||
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int neighflag,newton_pair;
|
||||
|
||||
@ -172,12 +172,12 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -254,8 +254,6 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev_all.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -266,6 +264,8 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -314,7 +314,7 @@ void PairTersoffMODKokkos<DeviceType>::operator()(TagPairTersoffMODComputeHalf<N
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
|
||||
F_FLOAT fi[3], fj[3], fk[3];
|
||||
|
||||
@ -331,7 +331,7 @@ void PairTersoffMODKokkos<DeviceType>::operator()(TagPairTersoffMODComputeHalf<N
|
||||
int j = d_neighbors_short(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtype = type(j);
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
if (itag > jtag) {
|
||||
if ((itag+jtag) % 2 == 0) continue;
|
||||
|
||||
@ -197,16 +197,16 @@ class PairTersoffMODKokkos : public PairTersoffMOD {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
|
||||
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
|
||||
typedef typename tdual_ffloat_2d_n7::t_host t_host_ffloat_2d_n7;
|
||||
|
||||
typename ArrayTypes<DeviceType>::t_neighbors_2d d_neighbors;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_ilist;
|
||||
typename ArrayTypes<DeviceType>::t_int_1d_randomread d_numneigh;
|
||||
typename AT::t_neighbors_2d d_neighbors;
|
||||
typename AT::t_int_1d_randomread d_ilist;
|
||||
typename AT::t_int_1d_randomread d_numneigh;
|
||||
//NeighListKokkos<DeviceType> k_list;
|
||||
|
||||
int neighflag,newton_pair;
|
||||
|
||||
@ -186,12 +186,12 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -268,8 +268,6 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev_all.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -280,6 +278,8 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
@ -328,7 +328,7 @@ void PairTersoffZBLKokkos<DeviceType>::operator()(TagPairTersoffZBLComputeHalf<N
|
||||
const X_FLOAT ytmp = x(i,1);
|
||||
const X_FLOAT ztmp = x(i,2);
|
||||
const int itype = type(i);
|
||||
const int itag = tag(i);
|
||||
const tagint itag = tag(i);
|
||||
|
||||
F_FLOAT fi[3], fj[3], fk[3];
|
||||
|
||||
@ -345,7 +345,7 @@ void PairTersoffZBLKokkos<DeviceType>::operator()(TagPairTersoffZBLComputeHalf<N
|
||||
int j = d_neighbors_short(i,jj);
|
||||
j &= NEIGHMASK;
|
||||
const int jtype = type(j);
|
||||
const int jtag = tag(j);
|
||||
const tagint jtag = tag(j);
|
||||
|
||||
if (itag > jtag) {
|
||||
if ((itag+jtag) % 2 == 0) continue;
|
||||
|
||||
@ -202,8 +202,8 @@ class PairTersoffZBLKokkos : public PairTersoffZBL {
|
||||
|
||||
DAT::tdual_efloat_1d k_eatom;
|
||||
DAT::tdual_virial_array k_vatom;
|
||||
DAT::t_efloat_1d d_eatom;
|
||||
DAT::t_virial_array d_vatom;
|
||||
typename ArrayTypes<DeviceType>::t_efloat_1d d_eatom;
|
||||
typename ArrayTypes<DeviceType>::t_virial_array d_vatom;
|
||||
|
||||
typedef Kokkos::DualView<F_FLOAT**[7],Kokkos::LayoutRight,DeviceType> tdual_ffloat_2d_n7;
|
||||
typedef typename tdual_ffloat_2d_n7::t_dev_const_randomread t_ffloat_2d_n7_randomread;
|
||||
|
||||
@ -12,8 +12,7 @@
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
/* ----------------------------------------------------------------------
|
||||
Contributing author: Stan Moore (SNL)
|
||||
Anders Hafreager (UiO), andershaf@gmail.com
|
||||
Contributing author: Anders Hafreager (UiO), andershaf@gmail.com
|
||||
------------------------------------------------------------------------- */
|
||||
|
||||
#include <math.h>
|
||||
@ -88,12 +87,12 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
if (eflag_atom) {
|
||||
memory->destroy_kokkos(k_eatom,eatom);
|
||||
memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom");
|
||||
d_eatom = k_eatom.d_view;
|
||||
d_eatom = k_eatom.view<DeviceType>();
|
||||
}
|
||||
if (vflag_atom) {
|
||||
memory->destroy_kokkos(k_vatom,vatom);
|
||||
memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom");
|
||||
d_vatom = k_vatom.d_view;
|
||||
d_vatom = k_vatom.view<DeviceType>();
|
||||
}
|
||||
|
||||
atomKK->sync(execution_space,datamask_read);
|
||||
@ -182,8 +181,6 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
virial[5] += ev_all.v[5];
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
if (eflag_atom) {
|
||||
k_eatom.template modify<DeviceType>();
|
||||
k_eatom.template sync<LMPHostType>();
|
||||
@ -194,6 +191,8 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
|
||||
k_vatom.template sync<LMPHostType>();
|
||||
}
|
||||
|
||||
if (vflag_fdotr) pair_virial_fdotr_compute(this);
|
||||
|
||||
copymode = 0;
|
||||
}
|
||||
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user