Merge pull request #3933 from lammps/maintenance

Second Set of Collected Bug Fixes and Maintenance Updates for 2 August 2023 Stable release
workaround hack for macOS
2023-12-14 21:09:30 -05:00 · 2023-12-14 18:07:50 -05:00 · 2023-12-14 17:32:20 -05:00 · 2023-12-14 16:48:25 -05:00 · 2023-12-14 16:24:22 -05:00 · 2023-12-14 16:00:10 -05:00
189 changed files with 2855 additions and 2059 deletions
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@ -151,12 +151,12 @@ tools/vim/*           @hammondkd
 unittest/*            @akohlmey

 # cmake
-cmake/*               @rbberger
+cmake/*               @akohlmey
 cmake/Modules/LAMMPSInterfacePlugin.cmake @akohlmey
 cmake/Modules/MPI4WIN.cmake @akohlmey
 cmake/Modules/OpenCLLoader.cmake @akohlmey
-cmake/Modules/Packages/COLVARS.cmake @rbberger @giacomofiorin
-cmake/Modules/Packages/KIM.cmake @rbberger @ellio167
+cmake/Modules/Packages/COLVARS.cmake @giacomofiorin
+cmake/Modules/Packages/KIM.cmake @ellio167
 cmake/presets/*.cmake @akohlmey

 # python
--- a/cmake/CMakeLists.txt
+++ b/cmake/CMakeLists.txt
@ -2,7 +2,6 @@
 ########################################
 # CMake build system
 # This file is part of LAMMPS
-# Created by Christoph Junghans and Richard Berger
 cmake_minimum_required(VERSION 3.10)
 ########################################
 # set policy to silence warnings about ignoring <PackageName>_ROOT but use it
@ -112,7 +111,7 @@ if(CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
    if(CMAKE_CXX_COMPILER_VERSION VERSION_EQUAL 17.3 OR CMAKE_CXX_COMPILER_VERSION VERSION_EQUAL 17.4)
      set(CMAKE_TUNE_DEFAULT "-xCOMMON-AVX512")
    else()
-      set(CMAKE_TUNE_DEFAULT "-xHost -fp-model fast=2 -no-prec-div -qoverride-limits -diag-disable=10441 -diag-disable=2196")
+      set(CMAKE_TUNE_DEFAULT "-xHost -fp-model fast=2 -no-prec-div -qoverride-limits -diag-disable=10441 -diag-disable=11074 -diag-disable=11076 -diag-disable=2196")
    endif()
  endif()
 endif()
@ -151,6 +150,7 @@ if(MSVC)
    add_compile_options(/Zc:__cplusplus)
    add_compile_options(/wd4244)
    add_compile_options(/wd4267)
+    add_compile_options(/wd4250)
    add_compile_options(/EHsc)
  endif()
  add_compile_definitions(_CRT_SECURE_NO_WARNINGS)
@ -438,6 +438,18 @@ if(BUILD_OMP)
  target_link_libraries(lmp PRIVATE OpenMP::OpenMP_CXX)
 endif()

+# lower C++ standard for fmtlib sources when using Intel classic compiler
+if((CMAKE_CXX_COMPILER_ID STREQUAL "Intel") AND (CMAKE_CXX_STANDARD GREATER_EQUAL 17)
+    AND (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 2021.10))
+  message(STATUS "Lowering C++ standard for compiling fmtlib sources with Intel Classic compiler")
+  get_filename_component(LMP_UTILS_SRC "${LAMMPS_SOURCE_DIR}/utils.cpp" ABSOLUTE)
+  get_filename_component(LMP_VARIABLE_SRC "${LAMMPS_SOURCE_DIR}/variable.cpp" ABSOLUTE)
+  get_filename_component(FMT_FORMAT_SRC "${LAMMPS_SOURCE_DIR}/fmtlib_format.cpp" ABSOLUTE)
+  get_filename_component(FMT_OS_SRC "${LAMMPS_SOURCE_DIR}/fmtlib_os.cpp" ABSOLUTE)
+  set_source_files_properties("${FMT_FORMAT_SRC}" "${FMT_OS_SRC}" "${LMP_VARIABLE_SRC}" "${LMP_UTILS_SRC}"
+          PROPERTIES COMPILE_OPTIONS "-std=c++14")
+endif()
+
 if(PKG_MSCG OR PKG_ATC OR PKG_AWPMD OR PKG_ML-QUIP OR PKG_ML-POD OR PKG_ELECTRODE OR BUILD_TOOLS)
  enable_language(C)
  if (NOT USE_INTERNAL_LINALG)
--- a/cmake/Modules/LAMMPSUtils.cmake
+++ b/cmake/Modules/LAMMPSUtils.cmake
@ -83,17 +83,17 @@ function(check_for_autogen_files source_dir)
    file(GLOB SRC_AUTOGEN_FILES ${CONFIGURE_DEPENDS} ${source_dir}/style_*.h)
    file(GLOB SRC_AUTOGEN_PACKAGES ${CONFIGURE_DEPENDS} ${source_dir}/packages_*.h)
    list(APPEND SRC_AUTOGEN_FILES ${SRC_AUTOGEN_PACKAGES} ${source_dir}/lmpinstalledpkgs.h ${source_dir}/lmpgitversion.h)
-    list(APPEND SRC_AUTOGEN_FILES ${SRC_AUTOGEN_PACKAGES} ${source_dir}/mliap_model_python_couple.h ${source_dir}/mliap_model_python_couple.cpp)
+    list(APPEND SRC_AUTOGEN_FILES ${source_dir}/mliap_model_python_couple.h ${source_dir}/mliap_model_python_couple.cpp)
    foreach(_SRC ${SRC_AUTOGEN_FILES})
      get_filename_component(FILENAME "${_SRC}" NAME)
      if(EXISTS ${source_dir}/${FILENAME})
        message(FATAL_ERROR "\n########################################################################\n"
-                              "Found header file(s) generated by the make-based build system\n"
-                              "\n"
-                              "Please run\n"
-                              "make -C ${source_dir} purge\n"
-                              "to remove\n"
-                              "########################################################################")
+                            "Found header file ${source_dir}/${FILENAME} generated by the make-based build system\n"
+                            "\n"
+                            "Please run\n"
+                            "make -C ${source_dir} purge\n"
+                            "to remove\n"
+                            "########################################################################")
      endif()
    endforeach()
 endfunction()
--- a/cmake/Modules/Packages/ML-PACE.cmake
+++ b/cmake/Modules/Packages/ML-PACE.cmake
@ -1,6 +1,6 @@
-set(PACELIB_URL "https://github.com/ICAMS/lammps-user-pace/archive/refs/tags/v.2023.01.3.fix.tar.gz" CACHE STRING "URL for PACE evaluator library sources")
+set(PACELIB_URL "https://github.com/ICAMS/lammps-user-pace/archive/refs/tags/v.2023.10.04.tar.gz" CACHE STRING "URL for PACE evaluator library sources")

-set(PACELIB_MD5 "4f0b3b5b14456fe9a73b447de3765caa" CACHE STRING "MD5 checksum of PACE evaluator library tarball")
+set(PACELIB_MD5 "70ff79f4e59af175e55d24f3243ad1ff" CACHE STRING "MD5 checksum of PACE evaluator library tarball")
 mark_as_advanced(PACELIB_URL)
 mark_as_advanced(PACELIB_MD5)
 GetFallbackURL(PACELIB_URL PACELIB_FALLBACK)
--- a/cmake/presets/kokkos-cuda.cmake
+++ b/cmake/presets/kokkos-cuda.cmake
@ -6,6 +6,8 @@ set(Kokkos_ENABLE_SERIAL ON CACHE BOOL "" FORCE)
 set(Kokkos_ENABLE_CUDA   ON CACHE BOOL "" FORCE)
 set(Kokkos_ARCH_PASCAL60 ON CACHE BOOL "" FORCE)
 set(BUILD_OMP ON CACHE BOOL "" FORCE)
+get_filename_component(NVCC_WRAPPER_CMD ${CMAKE_CURRENT_SOURCE_DIR}/../lib/kokkos/bin/nvcc_wrapper ABSOLUTE)
+set(CMAKE_CXX_COMPILER ${NVCC_WRAPPER_CMD} CACHE FILEPATH "" FORCE)

 # hide deprecation warnings temporarily for stable release
 set(Kokkos_ENABLE_DEPRECATION_WARNINGS OFF CACHE BOOL "" FORCE)
--- a/doc/Makefile
+++ b/doc/Makefile
@ -63,6 +63,7 @@ help:
 	@echo "  anchor_check  scan for duplicate anchor labels"
 	@echo "  style_check   check for complete and consistent style lists"
 	@echo "  package_check check for complete and consistent package lists"
+	@echo "  role_check    check for misformatted role keywords"
 	@echo "  spelling      spell-check the manual"

 # ------------------------------------------
@ -98,6 +99,7 @@ html: xmlgen $(VENV) $(SPHINXCONFIG)/conf.py $(ANCHORCHECK) $(MATHJAX)
 		env LC_ALL=C grep -n '[^ -~]' $(RSTDIR)/*.rst ;\
 		env LC_ALL=C grep -n ' :[a-z]\+`' $(RSTDIR)/*.rst ;\
 		env LC_ALL=C grep -n ' `[^`]\+<[a-z][^`]\+`[^_]' $(RSTDIR)/*.rst ;\
+		env LC_ALL=C grep -n ':\(ref\|doc\):[^`]' $(RSTDIR)/*.rst ;\
 		$(PYTHON) $(BUILDDIR)/utils/check-styles.py -s ../src -d src ;\
 		echo "############################################" ;\
 		deactivate ;\
@ -179,6 +181,7 @@ pdf: xmlgen $(VENV) $(SPHINXCONFIG)/conf.py $(ANCHORCHECK)
 		env LC_ALL=C grep -n '[^ -~]' $(RSTDIR)/*.rst ;\
 		env LC_ALL=C grep -n ' :[a-z]\+`' $(RSTDIR)/*.rst ;\
 		env LC_ALL=C grep -n ' `[^`]\+<[a-z][^`]\+`[^_]' $(RSTDIR)/*.rst ;\
+		env LC_ALL=C grep -n ':\(ref\|doc\):[^`]' $(RSTDIR)/*.rst ;\
 		$(PYTHON) utils/check-styles.py -s ../src -d src ;\
 		echo "############################################" ;\
 		deactivate ;\
@ -227,6 +230,7 @@ char_check :
 role_check :
 	@( env LC_ALL=C grep -n ' :[a-z]\+`' $(RSTDIR)/*.rst && exit 1 || : )
 	@( env LC_ALL=C grep -n ' `[^`]\+<[a-z][^`]\+`[^_]' $(RSTDIR)/*.rst && exit 1 || : )
+	@( env LC_ALL=C grep -n ':\(ref\|doc\):[^`]' $(RSTDIR)/*.rst && exit 1 || : )

 link_check : $(VENV) html
 	@(\
--- a/doc/src/Howto_lammps_gui.rst
+++ b/doc/src/Howto_lammps_gui.rst
@ -90,8 +90,8 @@ The run can be stopped cleanly by using either the ``Stop LAMMPS`` entry
 in the ``Run`` menu, the hotkey `Ctrl-/` (`Command-/` on macOS), or
 clicking on the red button in the status bar.  This will cause that the
 running LAMMPS process will complete the current iteration and then
-stop.  This is equivalent to the command `timer timeout 0 <timer>` and
-implemented by calling the :cpp:func:`lammps_force_timeout()` function
+stop.  This is equivalent to the command :doc:`timer timeout 0 <timer>`
+and implemented by calling the :cpp:func:`lammps_force_timeout()` function
 of the LAMMPS C-library interface.


--- a/doc/src/compute_saed.rst
+++ b/doc/src/compute_saed.rst
@ -68,7 +68,7 @@ reciprocal lattice nodes. The mesh spacing is defined either (a)  by
 the entire simulation domain or (b) manually using selected values as
 shown in the 2D diagram below.

-.. image:: img/saed_mesh.jpg
+.. image:: img/saed_mesh.png
   :scale: 75%
   :align: center

--- a/doc/src/compute_xrd.rst
+++ b/doc/src/compute_xrd.rst
@ -72,7 +72,7 @@ reciprocal lattice nodes. The mesh spacing is defined either (a) by the entire
 simulation domain or (b) manually using selected values as
 shown in the 2D diagram below.

-.. image:: img/xrd_mesh.jpg
+.. image:: img/xrd_mesh.png
   :scale: 75%
   :align: center

--- a/doc/src/fix_adapt_fep.rst
+++ b/doc/src/fix_adapt_fep.rst
@ -307,7 +307,9 @@ the :doc:`run <run>` command.  This fix is not invoked during

 Restrictions
 """"""""""""
- none
+
+The keyword "scale yes" is not supported for scaling per-atom parameters
+diameter and change. You can use :doc:`fix adapt <fix_adapt>` for those.

 Related commands
 """"""""""""""""
--- a/doc/src/fix_atom_swap.rst
+++ b/doc/src/fix_atom_swap.rst
@ -181,6 +181,12 @@ This fix is part of the MC package.  It is only enabled if LAMMPS was
 built with that package.  See the :doc:`Build package <Build_package>`
 doc page for more info.

+This fix cannot be used with systems that do not have per-type masses
+(e.g. atom style sphere) since the implemented algorithm pre-computes
+velocity rescaling factors from per-type masses and ignores any per-atom
+masses, if present.  In case both, per-type and per-atom masses are
+present, a warning is printed.
+
 Related commands
 """"""""""""""""

--- a/doc/src/fix_ave_chunk.rst
+++ b/doc/src/fix_ave_chunk.rst
@ -541,10 +541,10 @@ Restrictions
 Related commands
 """"""""""""""""

-:doc:`compute <compute>`, :doc:`fix ave/atom <fix_ave_atom>`, `fix
-:doc:ave/histo <fix_ave_histo>`, :doc:`fix ave/time <fix_ave_time>`,
-:doc:`variable <variable>`, :doc:`fix ave/correlate
-:doc:<fix_ave_correlate>`, `fix ave/atogrid <fix_ave_grid>`
+:doc:`compute <compute>`, :doc:`fix ave/atom <fix_ave_atom>`,
+:doc:`fix ave/histo <fix_ave_histo>`, :doc:`fix ave/time <fix_ave_time>`,
+:doc:`variable <variable>`, :doc:`fix ave/correlate <fix_ave_correlate>`,
+:doc:`fix ave/grid <fix_ave_grid>`


 Default
--- a/doc/src/fix_nh_uef.rst
+++ b/doc/src/fix_nh_uef.rst
@ -62,7 +62,7 @@ performed using the :doc:`fix deform <fix_deform>`, :doc:`fix nvt/sllod
 <fix_nvt_sllod>`, and :doc:`compute temp/deform <compute_temp_deform>`
 commands.

-The applied flow field is set by the *eps* keyword. The values
+The applied flow field is set by the *erate* keyword. The values
 *edot_x* and *edot_y* correspond to the strain rates in the xx and yy
 directions.  It is implicitly assumed that the flow field is
 traceless, and therefore the strain rate in the zz direction is eqal
--- a/doc/src/img/saed_mesh.jpg
+++ b/doc/src/img/saed_mesh.jpg
--- a/doc/src/img/saed_mesh.png
+++ b/doc/src/img/saed_mesh.png
--- a/doc/src/img/xrd_mesh.jpg
+++ b/doc/src/img/xrd_mesh.jpg
--- a/doc/src/img/xrd_mesh.png
+++ b/doc/src/img/xrd_mesh.png
--- a/doc/src/improper_amoeba.rst
+++ b/doc/src/improper_amoeba.rst
@ -68,8 +68,8 @@ for more info.
 Related commands
 """"""""""""""""

-:doc:`improper_coeff <improper_coeff>`, `improper_harmonic
-:doc:<improper_harmonic>`
+:doc:`improper_coeff <improper_coeff>`,
+:doc:`improper_harmonic <improper_harmonic>`

 Default
 """""""
--- a/doc/src/pair_beck.rst
+++ b/doc/src/pair_beck.rst
@ -30,11 +30,11 @@ Description

 Style *beck* computes interactions based on the potential by
 :ref:`(Beck) <Beck>`, originally designed for simulation of Helium.  It
-includes truncation at a cutoff distance Rc.
+includes truncation at a cutoff distance :math:`r_c`.

 .. math::

-   E(r) &= A \exp\left[-\alpha r - \beta r^6\right] - \frac{B}{\left(r^2+a^2\right)^3} \left(1+\frac{2.709+3a^2}{r^2+a^2}\right) \qquad r < R_c \\
+   E(r) &= A \exp\left[-\alpha r - \beta r^6\right] - \frac{B}{\left(r^2+a^2\right)^3} \left(1+\frac{2.709+3a^2}{r^2+a^2}\right) \qquad r < r_c \\

 The following coefficients must be defined for each pair of atoms
 types via the :doc:`pair_coeff <pair_coeff>` command as in the examples
@ -50,7 +50,7 @@ commands.
 * cutoff (distance units)

 The last coefficient is optional.  If not specified, the global cutoff
-:math:`R_c` is used.
+:math:`r_c` is used.

 ----------

--- a/doc/src/pair_lj_smooth_linear.rst
+++ b/doc/src/pair_lj_smooth_linear.rst
@ -31,13 +31,13 @@ Style *lj/smooth/linear* computes a truncated and force-shifted LJ
 interaction (aka Shifted Force Lennard-Jones) that combines the
 standard 12/6 Lennard-Jones function and subtracts a linear term based
 on the cutoff distance, so that both, the potential and the force, go
-continuously to zero at the cutoff Rc :ref:`(Toxvaerd) <Toxvaerd>`:
+continuously to zero at the cutoff :math:`r_c` :ref:`(Toxvaerd) <Toxvaerd>`:

 .. math::

   \phi\left(r\right) & =  4 \epsilon \left[ \left(\frac{\sigma}{r}\right)^{12} -
                       \left(\frac{\sigma}{r}\right)^6 \right] \\
-   E\left(r\right) & =  \phi\left(r\right)  - \phi\left(R_c\right) - \left(r - R_c\right) \left.\frac{d\phi}{d r} \right|_{r=R_c}       \qquad r < R_c
+   E\left(r\right) & =  \phi\left(r\right)  - \phi\left(r_c\right) - \left(r - r_c\right) \left.\frac{d\phi}{d r} \right|_{r=r_c}       \qquad r < r_c

 The following coefficients must be defined for each pair of atoms
 types via the :doc:`pair_coeff <pair_coeff>` command as in the examples
@ -77,8 +77,9 @@ tail option for adding long-range tail corrections to energy and
 pressure, since the energy of the pair interaction is smoothed to 0.0
 at the cutoff.

-This pair style writes its information to :doc:`binary restart files <restart>`, so pair_style and pair_coeff commands do not need
-to be specified in an input script that reads a restart file.
+This pair style writes its information to :doc:`binary restart files <restart>`,
+so pair_style and pair_coeff commands do not need to be specified
+in an input script that reads a restart file.

 This pair style can only be used via the *pair* keyword of the
 :doc:`run_style respa <run_style>` command.  It does not support the
--- a/doc/src/pair_mie.rst
+++ b/doc/src/pair_mie.rst
@ -35,7 +35,7 @@ The *mie/cut* style computes the Mie potential, given by
   E =  C \epsilon \left[ \left(\frac{\sigma}{r}\right)^{\gamma_{rep}} - \left(\frac{\sigma}{r}\right)^{\gamma_{att}} \right]
                         \qquad r < r_c

-Rc is the cutoff and C is a function that depends on the repulsive and
+:math:`r_c` is the cutoff and C is a function that depends on the repulsive and
 attractive exponents, given by:

 .. math::
--- a/doc/src/pair_morse.rst
+++ b/doc/src/pair_morse.rst
@ -53,7 +53,7 @@ Style *morse* computes pairwise interactions with the formula
   E = D_0 \left[ e^{- 2 \alpha (r - r_0)} - 2 e^{- \alpha (r - r_0)} \right]
       \qquad r < r_c

-Rc is the cutoff.
+:math:`r_c` is the cutoff.

 The following coefficients must be defined for each pair of atoms
 types via the :doc:`pair_coeff <pair_coeff>` command as in the examples
@ -78,7 +78,7 @@ so that both, potential energy and force, go to zero at the cut-off:
 .. math::

   \phi\left(r\right) & =  D_0 \left[ e^{- 2 \alpha (r - r_0)} - 2 e^{- \alpha (r - r_0)} \right] \qquad r < r_c \\
-   E\left(r\right) & =  \phi\left(r\right)  - \phi\left(R_c\right) - \left(r - R_c\right) \left.\frac{d\phi}{d r} \right|_{r=R_c}       \qquad r < R_c
+   E\left(r\right) & =  \phi\left(r\right)  - \phi\left(r_c\right) - \left(r - r_c\right) \left.\frac{d\phi}{d r} \right|_{r=r_c}       \qquad r < r_c

 The syntax of the pair_style and pair_coeff commands are the same for
 the *morse* and *morse/smooth/linear* styles.
--- a/doc/src/pair_soft.rst
+++ b/doc/src/pair_soft.rst
@ -44,8 +44,9 @@ It is useful for pushing apart overlapping atoms, since it does not
 blow up as r goes to 0.  A is a prefactor that can be made to vary in
 time from the start to the end of the run (see discussion below),
 e.g. to start with a very soft potential and slowly harden the
-interactions over time.  Rc is the cutoff.  See the :doc:`fix nve/limit <fix_nve_limit>` command for another way to push apart
-overlapping atoms.
+interactions over time.  :math:`r_c` is the cutoff.
+See the :doc:`fix nve/limit <fix_nve_limit>` command for another way
+to push apart overlapping atoms.

 The following coefficients must be defined for each pair of atom types
 via the :doc:`pair_coeff <pair_coeff>` command as in the examples above,
--- a/doc/src/pair_spica.rst
+++ b/doc/src/pair_spica.rst
@ -81,7 +81,7 @@ given by

 as required for the SPICA (formerly called SDK) and the pSPICA Coarse-grained MD parameterization discussed in
 :ref:`(Shinoda) <Shinoda3>`, :ref:`(DeVane) <DeVane>`, :ref:`(Seo) <Seo>`, and :ref:`(Miyazaki) <Miyazaki>`.
-Rc is the cutoff.
+:math:`r_c` is the cutoff.
 Summary information on these force fields can be found at https://www.spica-ff.org

 Style *lj/spica/coul/long* computes the adds Coulombic interactions
--- a/doc/utils/sphinx-config/LAMMPSLexer.py
+++ b/doc/utils/sphinx-config/LAMMPSLexer.py
@ -76,12 +76,12 @@ class LAMMPSLexer(RegexLexer):
            include('conditionals'),
            include('keywords'),
            (r'#.*?\n', Comment),
-            ('"', String, 'string'),
-            ('\'', String, 'single_quote_string'),
+            (r'"', String, 'string'),
+            (r'\'', String, 'single_quote_string'),
            (r'[0-9]+:[0-9]+(:[0-9]+)?', Number),
            (r'[0-9]+(\.[0-9]+)?([eE]\-?[0-9]+)?', Number),
-            ('\$?\(', Name.Variable, 'expression'),
-            ('\$\{', Name.Variable, 'variable'),
+            (r'\$?\(', Name.Variable, 'expression'),
+            (r'\$\{', Name.Variable, 'variable'),
            (r'[\w_\.\[\]]+', Name),
            (r'\$[\w_]+', Name.Variable),
            (r'\s+', Whitespace),
@ -97,21 +97,21 @@ class LAMMPSLexer(RegexLexer):
        ]
        ,
        'variable' : [
-            ('[^\}]+', Name.Variable),
-            ('\}', Name.Variable, '#pop'),
+            (r'[^\}]+', Name.Variable),
+            (r'\}', Name.Variable, '#pop'),
        ],
        'string' : [
-            ('[^"]+', String),
-            ('"', String, '#pop'),
+            (r'[^"]+', String),
+            (r'"', String, '#pop'),
        ],
        'single_quote_string' : [
-            ('[^\']+', String),
-            ('\'', String, '#pop'),
+            (r'[^\']+', String),
+            (r'\'', String, '#pop'),
        ],
        'expression' : [
-            ('[^\(\)]+', Name.Variable),
-            ('\(', Name.Variable, 'expression'),
-            ('\)', Name.Variable, '#pop'),
+            (r'[^\(\)]+', Name.Variable),
+            (r'\(', Name.Variable, 'expression'),
+            (r'\)', Name.Variable, '#pop'),
        ],
        'modify_cmd' : [
            (r'[\w_\-\.\[\]]+', Name.Variable.Identifier),
--- a/lib/pace/Install.py
+++ b/lib/pace/Install.py
@ -18,11 +18,11 @@ from install_helpers import fullpath, geturl, checkmd5sum, getfallback
 # settings

 thisdir = fullpath('.')
-version ='v.2023.01.3.fix'
+version ='v.2023.10.04'

 # known checksums for different PACE versions. used to validate the download.
 checksums = { \
-    'v.2023.01.3.fix': '4f0b3b5b14456fe9a73b447de3765caa'
+    'v.2023.10.04': '70ff79f4e59af175e55d24f3243ad1ff'
 }

 parser = ArgumentParser(prog='Install.py', description="LAMMPS library build wrapper script")
--- a/python/install.py
+++ b/python/install.py
@ -18,7 +18,7 @@ parser = ArgumentParser(prog='install.py',
                        description='LAMMPS python package installer script')

 parser.add_argument("-p", "--package", required=True,
-                    help="path to the LAMMPS Python package")
+                    help="path to the LAMMPS Python package folder")
 parser.add_argument("-l", "--lib", required=True,
                    help="path to the compiled LAMMPS shared library")
 parser.add_argument("-n", "--noinstall", action="store_true", default=False,
@ -34,15 +34,21 @@ args = parser.parse_args()

 if args.package:
  if not os.path.exists(args.package):
-    print("ERROR: LAMMPS package folder %s does not exist" % args.package)
+    print("\nERROR: LAMMPS package folder %s does not exist\n" % args.package)
    parser.print_help()
    sys.exit(1)
  else:
    args.package = os.path.abspath(args.package)
+    if ((os.path.basename(args.package) != "lammps")
+        and ((os.path.basename(os.path.dirname(args.package)) != "python"))):
+             print("\nERROR: LAMMPS package folder path %s does not end in %s\n"
+                   % (args.package, os.path.join("python", "lammps")))
+             parser.print_help()
+             sys.exit(1)

 if args.lib:
  if not os.path.exists(args.lib):
-    print("ERROR: LAMMPS shared library %s does not exist" % args.lib)
+    print("\nERROR: LAMMPS shared library %s does not exist\n" % args.lib)
    parser.print_help()
    sys.exit(1)
  else:
@ -50,7 +56,7 @@ if args.lib:

 if args.wheeldir:
  if not os.path.exists(args.wheeldir):
-    print("ERROR: directory %s to store the wheel does not exist" % args.wheeldir)
+    print("\nERROR: directory %s to store the wheel does not exist\n" % args.wheeldir)
    parser.print_help()
    sys.exit(1)
  else:
@ -58,7 +64,7 @@ if args.wheeldir:

 if args.versionfile:
  if not os.path.exists(args.versionfile):
-    print("ERROR: LAMMPS version file at %s does not exist" % args.versionfile)
+    print("\nERROR: LAMMPS version file at %s does not exist\n" % args.versionfile)
    parser.print_help()
    sys.exit(1)
  else:
--- a/python/lammps/core.py
+++ b/python/lammps/core.py
@ -379,8 +379,9 @@ class lammps(object):
          for i in range(narg):
            if type(cmdargs[i]) is str:
              cmdargs[i] = cmdargs[i].encode()
-          cargs = (c_char_p*narg)(*cmdargs)
-          self.lib.lammps_open.argtypes = [c_int, c_char_p*narg, MPI_Comm, c_void_p]
+          cargs = (c_char_p*(narg+1))(*cmdargs)
+          cargs[narg] = None
+          self.lib.lammps_open.argtypes = [c_int, c_char_p*(narg+1), MPI_Comm, c_void_p]
        else:
          self.lib.lammps_open.argtypes = [c_int, c_char_p, MPI_Comm, c_void_p]

@ -399,8 +400,9 @@ class lammps(object):
          for i in range(narg):
            if type(cmdargs[i]) is str:
              cmdargs[i] = cmdargs[i].encode()
-          cargs = (c_char_p*narg)(*cmdargs)
-          self.lib.lammps_open_no_mpi.argtypes = [c_int, c_char_p*narg, c_void_p]
+          cargs = (c_char_p*(narg+1))(*cmdargs)
+          cargs[narg] = None
+          self.lib.lammps_open_no_mpi.argtypes = [c_int, c_char_p*(narg+1), c_void_p]
          self.lmp = c_void_p(self.lib.lammps_open_no_mpi(narg,cargs,None))
        else:
          self.lib.lammps_open_no_mpi.argtypes = [c_int, c_char_p, c_void_p]
--- a/src/BOCS/fix_bocs.cpp
+++ b/src/BOCS/fix_bocs.cpp
@ -1024,7 +1024,10 @@ void FixBocs::final_integrate()

  if (pstat_flag) {
    if (pstyle == ISO) pressure->compute_scalar();
-    else pressure->compute_vector();
+    else {
+      temperature->compute_vector();
+      pressure->compute_vector();
+    }
    couple();
    pressure->addstep(update->ntimestep+1);
  }
@ -1961,6 +1964,7 @@ void FixBocs::nhc_press_integrate()
  int ich,i,pdof;
  double expfac,factor_etap,kecurrent;
  double kt = boltz * t_target;
+  double lkt_press;

  // Update masses, to preserve initial freq, if flag set

@ -2006,7 +2010,8 @@ void FixBocs::nhc_press_integrate()
    }
  }

-  double lkt_press = pdof * kt;
+  if (pstyle == ISO) lkt_press = kt;
+  else lkt_press = pdof * kt;
  etap_dotdot[0] = (kecurrent - lkt_press)/etap_mass[0];

  double ncfac = 1.0/nc_pchain;
--- a/src/DPD-MESO/fix_mvv_dpd.cpp
+++ b/src/DPD-MESO/fix_mvv_dpd.cpp
@ -63,13 +63,20 @@ int FixMvvDPD::setmask()
 void FixMvvDPD::init()
 {
  if (!atom->vest_flag)
-    error->all(FLERR,"Fix mvv/dpd requires atom attribute vest");
+    error->all(FLERR,"Fix mvv/dpd requires atom attribute vest e.g. from atom style mdpd");
+
+  if (!force->pair_match("^mdpd",0) && !force->pair_match("^dpd",0)) {
+    if (force->pair_match("^hybrid",0)) {
+      if (!(force->pair_match("^mdpd",0,1) || force->pair_match("^dpd",0),1)) {
+        error->all(FLERR, "Must use a dpd or mdpd pair style with fix mvv/dpd");
+      }
+    } else {
+      error->all(FLERR, "Must use a dpd or mdpd pair style with fix mvv/dpd");
+    }
+  }

  dtv = update->dt;
  dtf = 0.5 * update->dt * force->ftm2v;
-
-  if (!force->pair_match("^edpd",0) && !force->pair_match("^dpd",0))
-    error->all(FLERR, "Must use a dpd or edpd pair style with fix mvv/edpd");
 }

 /* ----------------------------------------------------------------------
--- a/src/DPD-MESO/fix_mvv_edpd.cpp
+++ b/src/DPD-MESO/fix_mvv_edpd.cpp
@ -71,11 +71,20 @@ int FixMvvEDPD::setmask()

 void FixMvvEDPD::init()
 {
+  if (!atom->edpd_flag) error->all(FLERR,"Fix mvv/edpd requires atom style edpd");
+
+  if (!force->pair_match("^edpd",0)) {
+    if (force->pair_match("^hybrid",0)) {
+      if (!force->pair_match("^edpd",0,1)) {
+        error->all(FLERR, "Must use pair style edpd with fix mvv/edpd");
+      }
+    } else {
+      error->all(FLERR, "Must use pair style edpd with fix mvv/edpd");
+    }
+  }
+
  dtv = update->dt;
  dtf = 0.5 * update->dt * force->ftm2v;
-
-  if (!force->pair_match("^edpd",0))
-    error->all(FLERR, "Must use pair style edpd with fix mvv/edpd");
 }

 /* ----------------------------------------------------------------------
--- a/src/DPD-MESO/fix_mvv_tdpd.cpp
+++ b/src/DPD-MESO/fix_mvv_tdpd.cpp
@ -69,10 +69,20 @@ int FixMvvTDPD::setmask()

 void FixMvvTDPD::init()
 {
+  if (!atom->tdpd_flag) error->all(FLERR,"Fix mvv/tdpd requires atom style tdpd");
+
+  if (!force->pair_match("^tdpd",0)) {
+    if (force->pair_match("^hybrid",0)) {
+      if (!force->pair_match("^tdpd",0,1)) {
+        error->all(FLERR, "Must use pair style tdpd with fix mvv/tdpd");
+      }
+    } else {
+      error->all(FLERR, "Must use pair style tdpd with fix mvv/tdpd");
+    }
+  }
+
  dtv = update->dt;
  dtf = 0.5 * update->dt * force->ftm2v;
-  if (!force->pair_match("^tdpd",0))
-    error->all(FLERR, "Must use pair style tdpd with fix mvv/tdpd");
 }

 /* ----------------------------------------------------------------------
--- a/src/Depend.sh
+++ b/src/Depend.sh
@ -64,6 +64,7 @@ fi

 if (test $1 = "COLLOID") then
  depend GPU
+  depend KOKKOS
  depend OPENMP
 fi

--- a/src/EXTRA-DUMP/dump_yaml.cpp
+++ b/src/EXTRA-DUMP/dump_yaml.cpp
@ -24,6 +24,8 @@

 using namespace LAMMPS_NS;

+static constexpr char special_chars[] = "{}[],&:*#?|-<>=!%@\\";
+
 /* ---------------------------------------------------------------------- */
 DumpYAML::DumpYAML(class LAMMPS *_lmp, int narg, char **args) :
    DumpCustom(_lmp, narg, args), thermo(false)
@ -67,7 +69,12 @@ void DumpYAML::write_header(bigint ndump)
      const auto &fields = th->get_fields();

      thermo_data += "thermo:\n  - keywords: [ ";
-      for (int i = 0; i < nfield; ++i) thermo_data += fmt::format("{}, ", keywords[i]);
+      for (int i = 0; i < nfield; ++i) {
+        if (keywords[i].find_first_of(special_chars) == std::string::npos)
+          thermo_data += fmt::format("{}, ", keywords[i]);
+        else
+          thermo_data += fmt::format("'{}', ", keywords[i]);
+      }
      thermo_data += "]\n  - data: [ ";

      for (int i = 0; i < nfield; ++i) {
@ -107,7 +114,12 @@ void DumpYAML::write_header(bigint ndump)
    if (domain->triclinic) fmt::print(fp, "  - [ {}, {}, {} ]\n", boxxy, boxxz, boxyz);

    fmt::print(fp, "keywords: [ ");
-    for (const auto &item : utils::split_words(columns)) fmt::print(fp, "{}, ", item);
+    for (const auto &item : utils::split_words(columns)) {
+      if (item.find_first_of(special_chars) == std::string::npos)
+        fmt::print(fp, "{}, ", item);
+      else
+        fmt::print(fp, "'{}', ", item);
+    }
    fputs(" ]\ndata:\n", fp);
  } else    // reset so that the remainder of the output is not multi-proc
    filewriter = 0;
--- a/src/FEP/fix_adapt_fep.cpp
+++ b/src/FEP/fix_adapt_fep.cpp
@ -40,17 +40,17 @@ using namespace LAMMPS_NS;
 using namespace FixConst;
 using namespace MathConst;

-enum{PAIR,KSPACE,ATOM};
-enum{DIAMETER,CHARGE};
+enum{PAIR, KSPACE, ATOM};
+enum{DIAMETER, CHARGE};

 /* ---------------------------------------------------------------------- */

 FixAdaptFEP::FixAdaptFEP(LAMMPS *lmp, int narg, char **arg) :
  Fix(lmp, narg, arg)
 {
-  if (narg < 5) error->all(FLERR,"Illegal fix adapt/fep command");
+  if (narg < 5) utils::missing_cmd_args(FLERR,"fix adapt/fep", error);
  nevery = utils::inumeric(FLERR,arg[3],false,lmp);
-  if (nevery < 0) error->all(FLERR,"Illegal fix adapt/fep command");
+  if (nevery < 0) error->all(FLERR,"Illegal fix adapt/fep every value {}", nevery);

  dynamic_group_allow = 1;
  create_attribute = 1;
@ -62,21 +62,21 @@ FixAdaptFEP::FixAdaptFEP(LAMMPS *lmp, int narg, char **arg) :
  int iarg = 4;
  while (iarg < narg) {
    if (strcmp(arg[iarg],"pair") == 0) {
-      if (iarg+6 > narg) error->all(FLERR,"Illegal fix adapt/fep command");
+      if (iarg+6 > narg) utils::missing_cmd_args(FLERR,"fix adapt/fep pair", error);
      nadapt++;
      iarg += 6;
    } else if (strcmp(arg[iarg],"kspace") == 0) {
-      if (iarg+2 > narg) error->all(FLERR,"Illegal fix adapt/fep command");
+      if (iarg+2 > narg) utils::missing_cmd_args(FLERR,"fix adapt/fep kspace", error);
      nadapt++;
      iarg += 2;
    } else if (strcmp(arg[iarg],"atom") == 0) {
-      if (iarg+4 > narg) error->all(FLERR,"Illegal fix adapt/fep command");
+      if (iarg+4 > narg) utils::missing_cmd_args(FLERR,"fix adapt/fep atom", error);
      nadapt++;
      iarg += 4;
    } else break;
  }

-  if (nadapt == 0) error->all(FLERR,"Illegal fix adapt/fep command");
+  if (nadapt == 0) error->all(FLERR,"Nothing to adapt in fix adapt/fep command");
  adapt = new Adapt[nadapt];

  // parse keywords
@ -136,11 +136,11 @@ FixAdaptFEP::FixAdaptFEP(LAMMPS *lmp, int narg, char **arg) :

  while (iarg < narg) {
    if (strcmp(arg[iarg],"reset") == 0) {
-      if (iarg+2 > narg) error->all(FLERR,"Illegal fix adapt/fep command");
+      if (iarg+2 > narg) utils::missing_cmd_args(FLERR,"fix adapt/fep reset", error);
      resetflag = utils::logical(FLERR,arg[iarg+1],false,lmp);
      iarg += 2;
    } else if (strcmp(arg[iarg],"scale") == 0) {
-      if (iarg+2 > narg) error->all(FLERR,"Illegal fix adapt/fep command");
+      if (iarg+2 > narg) utils::missing_cmd_args(FLERR,"fix adapt/fep scale", error);
      scaleflag = utils::logical(FLERR,arg[iarg+1],false,lmp);
      iarg += 2;
    } else if (strcmp(arg[iarg],"after") == 0) {
@ -165,21 +165,21 @@ FixAdaptFEP::FixAdaptFEP(LAMMPS *lmp, int narg, char **arg) :
 FixAdaptFEP::~FixAdaptFEP()
 {
  for (int m = 0; m < nadapt; m++) {
-    delete [] adapt[m].var;
+    delete[] adapt[m].var;
    if (adapt[m].which == PAIR) {
-      delete [] adapt[m].pstyle;
-      delete [] adapt[m].pparam;
+      delete[] adapt[m].pstyle;
+      delete[] adapt[m].pparam;
      memory->destroy(adapt[m].array_orig);
    }
  }
-  delete [] adapt;
+  delete[] adapt;

  // check nfix in case all fixes have already been deleted

  if (id_fix_diam && modify->nfix) modify->delete_fix(id_fix_diam);
  if (id_fix_chg && modify->nfix) modify->delete_fix(id_fix_chg);
-  delete [] id_fix_diam;
-  delete [] id_fix_chg;
+  delete[] id_fix_diam;
+  delete[] id_fix_chg;
 }

 /* ---------------------------------------------------------------------- */
@ -208,7 +208,7 @@ void FixAdaptFEP::post_constructor()
  id_fix_diam = nullptr;
  id_fix_chg = nullptr;

-  if (diam_flag) {
+  if (diam_flag && atom->radius_flag) {
    id_fix_diam = utils::strdup(id + std::string("_FIX_STORE_DIAM"));
    fix_diam = dynamic_cast<FixStoreAtom *>(
      modify->add_fix(fmt::format("{} {} STORE/ATOM 1 0 0 1", id_fix_diam,group->names[igroup])));
@ -226,7 +226,7 @@ void FixAdaptFEP::post_constructor()
    }
  }

-  if (chgflag) {
+  if (chgflag && atom->q_flag) {
    id_fix_chg = utils::strdup(id + std::string("_FIX_STORE_CHG"));
    fix_chg = dynamic_cast<FixStoreAtom *>(
      modify->add_fix(fmt::format("{} {} STORE/ATOM 1 0 0 1",id_fix_chg,group->names[igroup])));
@ -267,9 +267,9 @@ void FixAdaptFEP::init()

    ad->ivar = input->variable->find(ad->var);
    if (ad->ivar < 0)
-      error->all(FLERR,"Variable name for fix adapt/fep does not exist");
+      error->all(FLERR,"Variable name {} for fix adapt/fep does not exist", ad->var);
    if (!input->variable->equalstyle(ad->ivar))
-      error->all(FLERR,"Variable for fix adapt/fep is invalid style");
+      error->all(FLERR,"Variable {} for fix adapt/fep is invalid style", ad->var);

    if (ad->which == PAIR) {
      anypair = 1;
@ -285,8 +285,9 @@ void FixAdaptFEP::init()
      if (ptr == nullptr)
        error->all(FLERR,"Fix adapt/fep pair style param not supported");

-      ad->pdim = 2;
-      if (ad->pdim == 0) ad->scalar = (double *) ptr;
+      if (ad->pdim != 2)
+        error->all(FLERR,"Pair style parameter {} is not compatible with fix adapt/fep", ad->pparam);
+
      if (ad->pdim == 2) ad->array = (double **) ptr;

      // if pair hybrid, test that ilo,ihi,jlo,jhi are valid for sub-style
@ -434,6 +435,8 @@ void FixAdaptFEP::change_settings()

    } else if (ad->which == ATOM) {

+      if (scaleflag)
+        error->all(FLERR, "Keyword 'scale yes' is not supported with fix adapt/fep for 'atom'");
      // reset radius from diameter
      // also scale rmass to new value

--- a/src/KOKKOS/atom_kokkos.cpp
+++ b/src/KOKKOS/atom_kokkos.cpp
@ -44,6 +44,9 @@ AtomKokkos::AtomKokkos(LAMMPS *lmp) : Atom(lmp)

  h_tag_min = Kokkos::subview(h_tag_min_max,0);
  h_tag_max = Kokkos::subview(h_tag_min_max,1);
+
+  nprop_atom = 0;
+  fix_prop_atom = nullptr;
 }

 /* ---------------------------------------------------------------------- */
@ -112,6 +115,7 @@ AtomKokkos::~AtomKokkos()

  memoryKK->destroy_kokkos(k_dvector, dvector);
  dvector = nullptr;
+  delete [] fix_prop_atom;
 }

 /* ---------------------------------------------------------------------- */
@ -125,11 +129,37 @@ void AtomKokkos::init()

 /* ---------------------------------------------------------------------- */

+void AtomKokkos::update_property_atom()
+{
+  nprop_atom = 0;
+  std::vector<Fix *> prop_atom_fixes;
+  for (auto &ifix : modify->get_fix_by_style("^property/atom")) {
+    if (!ifix->kokkosable)
+      error->all(FLERR, "KOKKOS package requires a Kokkos-enabled version of fix property/atom");
+
+    ++nprop_atom;
+    prop_atom_fixes.push_back(ifix);
+  }
+
+  delete[] fix_prop_atom;
+  fix_prop_atom = new FixPropertyAtomKokkos *[nprop_atom];
+
+  int n = 0;
+  for (auto &ifix : prop_atom_fixes)
+    fix_prop_atom[n++] = dynamic_cast<FixPropertyAtomKokkos *>(ifix);
+}
+
+/* ---------------------------------------------------------------------- */
+
 void AtomKokkos::sync(const ExecutionSpace space, unsigned int mask)
 {
-  if (space == Device && lmp->kokkos->auto_sync) avecKK->modified(Host, mask);
+  if (space == Device && lmp->kokkos->auto_sync) {
+    avecKK->modified(Host, mask);
+    for (int n = 0; n < nprop_atom; n++) fix_prop_atom[n]->modified(Host, mask);
+  }

  avecKK->sync(space, mask);
+  for (int n = 0; n < nprop_atom; n++) fix_prop_atom[n]->sync(space, mask);
 }

 /* ---------------------------------------------------------------------- */
@ -137,13 +167,20 @@ void AtomKokkos::sync(const ExecutionSpace space, unsigned int mask)
 void AtomKokkos::modified(const ExecutionSpace space, unsigned int mask)
 {
  avecKK->modified(space, mask);
+  for (int n = 0; n < nprop_atom; n++) fix_prop_atom[n]->modified(space, mask);

-  if (space == Device && lmp->kokkos->auto_sync) avecKK->sync(Host, mask);
+  if (space == Device && lmp->kokkos->auto_sync) {
+    avecKK->sync(Host, mask);
+    for (int n = 0; n < nprop_atom; n++) fix_prop_atom[n]->sync(Host, mask);
+  }
 }

+/* ---------------------------------------------------------------------- */
+
 void AtomKokkos::sync_overlapping_device(const ExecutionSpace space, unsigned int mask)
 {
  avecKK->sync_overlapping_device(space, mask);
+  for (int n = 0; n < nprop_atom; n++) fix_prop_atom[n]->sync_overlapping_device(space, mask);
 }
 /* ---------------------------------------------------------------------- */

@ -375,7 +412,7 @@ AtomVec *AtomKokkos::new_avec(const std::string &style, int trysuffix, int &sfla
  int hybrid_substyle_flag = (avec != nullptr);

  AtomVec *avec = Atom::new_avec(style, trysuffix, sflag);
-  if (!avec->kokkosable) error->all(FLERR, "KOKKOS package requires a kokkos enabled atom_style");
+  if (!avec->kokkosable) error->all(FLERR, "KOKKOS package requires a Kokkos-enabled atom_style");

  if (!hybrid_substyle_flag)
    avecKK = dynamic_cast<AtomVecKokkos*>(avec);
--- a/src/KOKKOS/atom_kokkos.h
+++ b/src/KOKKOS/atom_kokkos.h
@ -14,6 +14,7 @@

 #include "atom.h"               // IWYU pragma: export
 #include "kokkos_type.h"
+#include "fix_property_atom_kokkos.h"

 #include <Kokkos_Sort.hpp>

@ -25,6 +26,8 @@ namespace LAMMPS_NS {
 class AtomKokkos : public Atom {
 public:
  bool sort_classic;
+  int nprop_atom;
+  FixPropertyAtomKokkos** fix_prop_atom;

  DAT::tdual_tagint_1d k_tag;
  DAT::tdual_int_1d k_type, k_mask;
@ -144,6 +147,7 @@ class AtomKokkos : public Atom {
  }

  void init() override;
+  void update_property_atom();
  void allocate_type_arrays() override;
  void sync(const ExecutionSpace space, unsigned int mask);
  void modified(const ExecutionSpace space, unsigned int mask);
--- a/src/KOKKOS/atom_vec_dpd_kokkos.cpp
+++ b/src/KOKKOS/atom_vec_dpd_kokkos.cpp
@ -963,7 +963,6 @@ void AtomVecDPDKokkos::sync(ExecutionSpace space, unsigned int mask)
    if (mask & UCG_MASK) atomKK->k_uCG.sync<LMPDeviceType>();
    if (mask & UCGNEW_MASK) atomKK->k_uCGnew.sync<LMPDeviceType>();
    if (mask & DUCHEM_MASK) atomKK->k_duChem.sync<LMPDeviceType>();
-    if (mask & DVECTOR_MASK) atomKK->k_dvector.sync<LMPDeviceType>();
  } else {
    if (mask & X_MASK) atomKK->k_x.sync<LMPHostType>();
    if (mask & V_MASK) atomKK->k_v.sync<LMPHostType>();
@ -980,7 +979,6 @@ void AtomVecDPDKokkos::sync(ExecutionSpace space, unsigned int mask)
    if (mask & UCG_MASK) atomKK->k_uCG.sync<LMPHostType>();
    if (mask & UCGNEW_MASK) atomKK->k_uCGnew.sync<LMPHostType>();
    if (mask & DUCHEM_MASK) atomKK->k_duChem.sync<LMPHostType>();
-    if (mask & DVECTOR_MASK) atomKK->k_dvector.sync<LMPHostType>();
  }
 }

@ -1019,8 +1017,6 @@ void AtomVecDPDKokkos::sync_overlapping_device(ExecutionSpace space, unsigned in
      perform_async_copy<DAT::tdual_efloat_1d>(atomKK->k_uCGnew,space);
    if ((mask & DUCHEM_MASK) && atomKK->k_duChem.need_sync<LMPDeviceType>())
      perform_async_copy<DAT::tdual_efloat_1d>(atomKK->k_duChem,space);
-    if ((mask & DVECTOR_MASK) && atomKK->k_dvector.need_sync<LMPDeviceType>())
-      perform_async_copy<DAT::tdual_float_2d>(atomKK->k_dvector,space);
  } else {
    if ((mask & X_MASK) && atomKK->k_x.need_sync<LMPHostType>())
      perform_async_copy<DAT::tdual_x_array>(atomKK->k_x,space);
@ -1052,8 +1048,6 @@ void AtomVecDPDKokkos::sync_overlapping_device(ExecutionSpace space, unsigned in
      perform_async_copy<DAT::tdual_efloat_1d>(atomKK->k_uCGnew,space);
    if ((mask & DUCHEM_MASK) && atomKK->k_duChem.need_sync<LMPHostType>())
      perform_async_copy<DAT::tdual_efloat_1d>(atomKK->k_duChem,space);
-    if ((mask & DVECTOR_MASK) && atomKK->k_dvector.need_sync<LMPHostType>())
-      perform_async_copy<DAT::tdual_float_2d>(atomKK->k_dvector,space);
  }
 }

@ -1077,7 +1071,6 @@ void AtomVecDPDKokkos::modified(ExecutionSpace space, unsigned int mask)
    if (mask & UCG_MASK) atomKK->k_uCG.modify<LMPDeviceType>();
    if (mask & UCGNEW_MASK) atomKK->k_uCGnew.modify<LMPDeviceType>();
    if (mask & DUCHEM_MASK) atomKK->k_duChem.modify<LMPDeviceType>();
-    if (mask & DVECTOR_MASK) atomKK->k_dvector.modify<LMPDeviceType>();
  } else {
    if (mask & X_MASK) atomKK->k_x.modify<LMPHostType>();
    if (mask & V_MASK) atomKK->k_v.modify<LMPHostType>();
@ -1094,6 +1087,5 @@ void AtomVecDPDKokkos::modified(ExecutionSpace space, unsigned int mask)
    if (mask & UCG_MASK) atomKK->k_uCG.modify<LMPHostType>();
    if (mask & UCGNEW_MASK) atomKK->k_uCGnew.modify<LMPHostType>();
    if (mask & DUCHEM_MASK) atomKK->k_duChem.modify<LMPHostType>();
-    if (mask & DVECTOR_MASK) atomKK->k_dvector.modify<LMPHostType>();
  }
 }
--- a/src/KOKKOS/atom_vec_kokkos.h
+++ b/src/KOKKOS/atom_vec_kokkos.h
@ -139,6 +139,8 @@ class AtomVecKokkos : virtual public AtomVec {

  DAT::tdual_int_1d k_count;

+ public:
+
  #ifdef LMP_KOKKOS_GPU
  template<class ViewType>
  Kokkos::View<typename ViewType::data_type,
--- a/src/KOKKOS/atom_vec_spin_kokkos.cpp
+++ b/src/KOKKOS/atom_vec_spin_kokkos.cpp
@ -586,6 +586,38 @@ int AtomVecSpinKokkos::unpack_exchange_kokkos(DAT::tdual_xfloat_2d &k_buf, int n
  }
 }

+/* ----------------------------------------------------------------------
+   clear extra forces starting at atom N
+   nbytes = # of bytes to clear for a per-atom vector
+   include f b/c this is invoked from within SPIN pair styles
+------------------------------------------------------------------------- */
+
+void AtomVecSpinKokkos::force_clear(int n, size_t nbytes)
+{
+  int nzero = (double)nbytes/sizeof(double);
+
+  if (nzero) {
+    atomKK->k_fm.clear_sync_state(); // will be cleared below
+    atomKK->k_fm_long.clear_sync_state(); // will be cleared below
+
+    // local variables for lambda capture
+
+    auto l_fm = atomKK->k_fm.d_view;
+    auto l_fm_long = atomKK->k_fm_long.d_view;
+
+    Kokkos::parallel_for(nzero, LAMMPS_LAMBDA(int i) {
+      l_fm(i,0) = 0.0;
+      l_fm(i,1) = 0.0;
+      l_fm(i,2) = 0.0;
+      l_fm_long(i,0) = 0.0;
+      l_fm_long(i,1) = 0.0;
+      l_fm_long(i,2) = 0.0;
+    });
+
+    atomKK->modified(Device,FM_MASK|FML_MASK);
+  }
+}
+
 /* ---------------------------------------------------------------------- */

 void AtomVecSpinKokkos::sync(ExecutionSpace space, unsigned int mask)
--- a/src/KOKKOS/atom_vec_spin_kokkos.h
+++ b/src/KOKKOS/atom_vec_spin_kokkos.h
@ -34,6 +34,7 @@ class AtomVecSpinKokkos : public AtomVecKokkos, public AtomVecSpin {
  AtomVecSpinKokkos(class LAMMPS *);
  void grow(int) override;
  void grow_pointers() override;
+  void force_clear(int, size_t) override;
  void sort_kokkos(Kokkos::BinSort<KeyViewType, BinOp> &Sorter) override;
  int pack_border_kokkos(int n, DAT::tdual_int_2d k_sendlist,
                         DAT::tdual_xfloat_2d buf,int iswap,
--- a/src/KOKKOS/fix_dt_reset_kokkos.cpp
+++ b/src/KOKKOS/fix_dt_reset_kokkos.cpp
@ -52,14 +52,6 @@ void FixDtResetKokkos<DeviceType>::init()
 {
  FixDtReset::init();

-  k_emax = Kokkos::DualView<double*, Kokkos::LayoutRight, DeviceType>("FixDtResetKokkos:gamma", 1);
-
-  k_emax.h_view(0) = emax;
-
-
-  k_emax.template modify<LMPHostType>();
-  k_emax.template sync<DeviceType>();
-
  if (utils::strmatch(update->integrate_style,"^respa"))
    error->all(FLERR,"Cannot (yet) use respa with Kokkos");
 }
@ -93,43 +85,38 @@ void FixDtResetKokkos<DeviceType>::end_of_step()

  MPI_Allreduce(MPI_IN_PLACE, &dt, 1, MPI_DOUBLE, MPI_MIN, world);

-  atomKK->modified(execution_space, F_MASK);
+  if (minbound) dt = MAX(dt, tmin);
+  if (maxbound) dt = MIN(dt, tmax);

-   if (minbound) dt = MAX(dt, tmin);
-   if (maxbound) dt = MIN(dt, tmax);
+  // if timestep didn't change, just return
+  // else reset update->dt and other classes that depend on it
+  // rRESPA, pair style, fixes

-   // if timestep didn't change, just return
-   // else reset update->dt and other classes that depend on it
-   // rRESPA, pair style, fixes
+  if (dt == update->dt) return;

-   if (dt == update->dt) return;
+  laststep = update->ntimestep;

-   laststep = update->ntimestep;
-
-   // calls to other classes that need to know timestep size changed
-   // similar logic is in Input::timestep()
-
-   update->update_time();
-   update->dt = dt;
-   update->dt_default = 0;
-   if (force->pair) force->pair->reset_dt();
-   for (int i = 0; i < modify->nfix; i++) modify->fix[i]->reset_dt();
-   output->reset_dt();
+  // calls to other classes that need to know timestep size changed
+  // similar logic is in Input::timestep()

+  update->update_time();
+  update->dt = dt;
+  update->dt_default = 0;
+  if (force->pair) force->pair->reset_dt();
+  for (auto &ifix : modify->get_fix_list()) ifix->reset_dt();
+  output->reset_dt();
 }

 /* ---------------------------------------------------------------------- */

 template<class DeviceType>
 KOKKOS_INLINE_FUNCTION
-void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetMass, const int &i, double &k_dt) const {
+void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetMass, const int &i, double &dt_min) const {

-  double dtv, dtf, dte, dtsq;
+  double dt, dtv, dtf, dte, dtsq;
  double vsq, fsq, massinv;
  double delx, dely, delz, delr;

-  double emax = k_emax.d_view(0);
-
  if (mask[i] & groupbit) {

    massinv = 1.0 / mass[type[i]];
@ -138,32 +125,31 @@ void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetMass, const int &i, d
    dtv = dtf = dte = BIG;
    if (vsq > 0.0) dtv = xmax / sqrt(vsq);
    if (fsq > 0.0) dtf = sqrt(2.0 * xmax / (ftm2v * sqrt(fsq) * massinv));
-    k_dt = MIN(dtv, dtf);
+    dt = MIN(dtv, dtf);
    if ((emax > 0.0) && (fsq * vsq > 0.0)) {
      dte = emax / sqrt(fsq * vsq) / sqrt(ftm2v * mvv2e);
-      k_dt = MIN(dt, dte);
+      dt = MIN(dt, dte);
    }
-    dtsq = k_dt * k_dt;
-    delx = k_dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v;
-    dely = k_dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v;
-    delz = k_dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v;
+    dtsq = dt * dt;
+    delx = dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v;
+    dely = dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v;
+    delz = dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v;
    delr = sqrt(delx * delx + dely * dely + delz * delz);
-    if (delr > xmax) k_dt *= xmax / delr;
+    if (delr > xmax) dt *= xmax / delr;
+    dt_min = MIN(dt_min,dt);
  }
- }
+}

 /* ---------------------------------------------------------------------- */

 template<class DeviceType>
 KOKKOS_INLINE_FUNCTION
-void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetRMass, const int &i, double &k_dt) const {
+void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetRMass, const int &i, double &dt_min) const {

-  double dtv, dtf, dte, dtsq;
+  double dt, dtv, dtf, dte, dtsq;
  double vsq, fsq, massinv;
  double delx, dely, delz, delr;

-  double emax = k_emax.d_view(0);
-
  if (mask[i] & groupbit) {

    massinv = 1.0 / rmass[i];
@ -172,17 +158,18 @@ void FixDtResetKokkos<DeviceType>::operator()(TagFixDtResetRMass, const int &i,
    dtv = dtf = dte = BIG;
    if (vsq > 0.0) dtv = xmax / sqrt(vsq);
    if (fsq > 0.0) dtf = sqrt(2.0 * xmax / (ftm2v * sqrt(fsq) * massinv));
-    k_dt = MIN(dtv, dtf);
+    dt = MIN(dtv, dtf);
    if ((emax > 0.0) && (fsq * vsq > 0.0)) {
      dte = emax / sqrt(fsq * vsq) / sqrt(ftm2v * mvv2e);
-      k_dt = MIN(dt, dte);
+      dt = MIN(dt, dte);
    }
-    dtsq = k_dt * k_dt;
-    delx = k_dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v;
-    dely = k_dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v;
-    delz = k_dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v;
+    dtsq = dt * dt;
+    delx = dt * v(i,0) + 0.5 * dtsq * massinv * f(i,0) * ftm2v;
+    dely = dt * v(i,1) + 0.5 * dtsq * massinv * f(i,1) * ftm2v;
+    delz = dt * v(i,2) + 0.5 * dtsq * massinv * f(i,2) * ftm2v;
    delr = sqrt(delx * delx + dely * dely + delz * delz);
-    if (delr > xmax) k_dt *= xmax / delr;
+    if (delr > xmax) dt *= xmax / delr;
+    dt_min = MIN(dt_min,dt);
  }
 }

--- a/src/KOKKOS/fix_enforce2d_kokkos.cpp
+++ b/src/KOKKOS/fix_enforce2d_kokkos.cpp
@ -17,15 +17,14 @@
 ------------------------------------------------------------------------- */

 #include "fix_enforce2d_kokkos.h"
+
 #include "atom_masks.h"
 #include "atom_kokkos.h"
 #include "comm.h"
 #include "error.h"

-
 using namespace LAMMPS_NS;

-
 template <class DeviceType>
 FixEnforce2DKokkos<DeviceType>::FixEnforce2DKokkos(LAMMPS *lmp, int narg, char **arg) :
  FixEnforce2D(lmp, narg, arg)
@ -34,21 +33,16 @@ FixEnforce2DKokkos<DeviceType>::FixEnforce2DKokkos(LAMMPS *lmp, int narg, char *
  atomKK = (AtomKokkos *) atom;
  execution_space = ExecutionSpaceFromDevice<DeviceType>::space;

-  datamask_read   = V_MASK | F_MASK | OMEGA_MASK | MASK_MASK
-          | TORQUE_MASK | ANGMOM_MASK;
-
-  datamask_modify = V_MASK | F_MASK | OMEGA_MASK
-          | TORQUE_MASK | ANGMOM_MASK;
+  datamask_read   = V_MASK | F_MASK | OMEGA_MASK | MASK_MASK | TORQUE_MASK | ANGMOM_MASK;
+  datamask_modify = V_MASK | F_MASK | OMEGA_MASK | TORQUE_MASK | ANGMOM_MASK;
 }

-
 template <class DeviceType>
 void FixEnforce2DKokkos<DeviceType>::setup(int vflag)
 {
  post_force(vflag);
 }

-
 template <class DeviceType>
 void FixEnforce2DKokkos<DeviceType>::post_force(int /*vflag*/)
 {
@ -66,7 +60,6 @@ void FixEnforce2DKokkos<DeviceType>::post_force(int /*vflag*/)
  if (atomKK->torque_flag)
    torque = atomKK->k_torque.view<DeviceType>();

-
  mask = atomKK->k_mask.view<DeviceType>();

  int nlocal = atomKK->nlocal;
@ -125,13 +118,6 @@ void FixEnforce2DKokkos<DeviceType>::post_force(int /*vflag*/)
  copymode = 0;

  atomKK->modified(execution_space,datamask_modify);
-
-  for (int m = 0; m < nfixlist; m++) {
-    atomKK->sync(flist[m]->execution_space,flist[m]->datamask_read);
-    flist[m]->enforce2d();
-    atomKK->modified(flist[m]->execution_space,flist[m]->datamask_modify);
-  }
-
 }


--- a/src/KOKKOS/fix_property_atom_kokkos.cpp
+++ b/src/KOKKOS/fix_property_atom_kokkos.cpp
@ -30,7 +30,46 @@ FixPropertyAtomKokkos::FixPropertyAtomKokkos(LAMMPS *lmp, int narg, char **arg)
  FixPropertyAtom(lmp, narg, arg)
 {
  atomKK = (AtomKokkos *) atom;
-  grow_arrays(atom->nmax);
+  kokkosable = 1;
+
+  dvector_flag = 0;
+  for (int nv = 0; nv < nvalue; nv++)
+    if (styles[nv] == DVEC) dvector_flag = 1;
+}
+
+/* ---------------------------------------------------------------------- */
+
+void FixPropertyAtomKokkos::post_constructor()
+{
+  atomKK->update_property_atom();
+
+  FixPropertyAtom::post_constructor();
+}
+
+/* ---------------------------------------------------------------------- */
+
+FixPropertyAtomKokkos::~FixPropertyAtomKokkos()
+{
+  // deallocate per-atom vectors in Atom class
+  // set ptrs to a null pointer, so they no longer exist for Atom class
+
+  for (int nv = 0; nv < nvalue; nv++) {
+    if (styles[nv] == MOLECULE) {
+      atom->molecule_flag = 0;
+      memoryKK->destroy_kokkos(atomKK->k_molecule,atom->molecule);
+      atom->molecule = nullptr;
+    } else if (styles[nv] == CHARGE) {
+      atom->q_flag = 0;
+      memoryKK->destroy_kokkos(atomKK->k_q,atom->q);
+      atom->q = nullptr;
+    } else if (styles[nv] == RMASS) {
+      atom->rmass_flag = 0;
+      memoryKK->destroy_kokkos(atomKK->k_rmass,atom->rmass);
+      atom->rmass = nullptr;
+    }
+  }
+
+  atomKK->update_property_atom();
 }

 /* ----------------------------------------------------------------------
@ -44,17 +83,17 @@ void FixPropertyAtomKokkos::grow_arrays(int nmax)
 {
  for (int nv = 0; nv < nvalue; nv++) {
    if (styles[nv] == MOLECULE) {
-      memory->grow(atom->molecule,nmax,"atom:molecule");
-      size_t nbytes = (nmax-nmax_old) * sizeof(tagint);
-      memset(&atom->molecule[nmax_old],0,nbytes);
+      atomKK->sync(Device,MOLECULE_MASK);
+      memoryKK->grow_kokkos(atomKK->k_molecule,atom->molecule,nmax,"atom:molecule");
+      atomKK->modified(Device,MOLECULE_MASK);
    } else if (styles[nv] == CHARGE) {
-      memory->grow(atom->q,nmax,"atom:q");
-      size_t nbytes = (nmax-nmax_old) * sizeof(double);
-      memset(&atom->q[nmax_old],0,nbytes);
+      atomKK->sync(Device,Q_MASK);
+      memoryKK->grow_kokkos(atomKK->k_q,atom->q,nmax,"atom:q");
+      atomKK->modified(Device,Q_MASK);
    } else if (styles[nv] == RMASS) {
-      memory->grow(atom->rmass,nmax,"atom:rmass");
-      size_t nbytes = (nmax-nmax_old) * sizeof(double);
-      memset(&atom->rmass[nmax_old],0,nbytes);
+      atomKK->sync(Device,RMASS_MASK);
+      memoryKK->grow_kokkos(atomKK->k_rmass,atom->rmass,nmax,"atom:rmass");
+      atomKK->modified(Device,RMASS_MASK);
    } else if (styles[nv] == TEMPERATURE) {
      memory->grow(atom->temperature, nmax, "atom:temperature");
      size_t nbytes = (nmax - nmax_old) * sizeof(double);
@ -69,7 +108,7 @@ void FixPropertyAtomKokkos::grow_arrays(int nmax)
      memset(&atom->ivector[index[nv]][nmax_old],0,nbytes);
    } else if (styles[nv] == DVEC) {
      atomKK->sync(Device,DVECTOR_MASK);
-      memoryKK->grow_kokkos(atomKK->k_dvector,atomKK->dvector,atomKK->k_dvector.extent(0),nmax,
+      memoryKK->grow_kokkos(atomKK->k_dvector,atom->dvector,atomKK->k_dvector.extent(0),nmax,
                          "atom:dvector");
      atomKK->modified(Device,DVECTOR_MASK);
    } else if (styles[nv] == IARRAY) {
@ -84,3 +123,62 @@ void FixPropertyAtomKokkos::grow_arrays(int nmax)
  }
  nmax_old = nmax;
 }
+
+/* ---------------------------------------------------------------------- */
+
+void FixPropertyAtomKokkos::sync(ExecutionSpace space, unsigned int mask)
+{
+  if (space == Device) {
+    if (molecule_flag && (mask & MOLECULE_MASK)) atomKK->k_molecule.sync<LMPDeviceType>();
+    if (q_flag && (mask & Q_MASK)) atomKK->k_q.sync<LMPDeviceType>();
+    if (rmass_flag && (mask & RMASS_MASK)) {atomKK->k_rmass.sync<LMPDeviceType>();}
+    if (dvector_flag && (mask & DVECTOR_MASK)) atomKK->k_dvector.sync<LMPDeviceType>();
+  } else {
+    if (molecule_flag && (mask & MOLECULE_MASK)) atomKK->k_molecule.sync<LMPHostType>();
+    if (q_flag && (mask & Q_MASK)) atomKK->k_q.sync<LMPHostType>();
+    if (rmass_flag && (mask & RMASS_MASK)) atomKK->k_rmass.sync<LMPHostType>();
+    if (dvector_flag && (mask & DVECTOR_MASK)) atomKK->k_dvector.sync<LMPHostType>();
+  }
+}
+
+/* ---------------------------------------------------------------------- */
+
+void FixPropertyAtomKokkos::sync_overlapping_device(ExecutionSpace space, unsigned int mask)
+{
+  if (space == Device) {
+    if ((mask & MOLECULE_MASK) && atomKK->k_molecule.need_sync<LMPDeviceType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_tagint_1d>(atomKK->k_molecule,space);
+    if ((mask & Q_MASK) && atomKK->k_q.need_sync<LMPDeviceType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_1d>(atomKK->k_q,space);
+    if ((mask & RMASS_MASK) && atomKK->k_rmass.need_sync<LMPDeviceType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_1d>(atomKK->k_rmass,space);
+    if ((mask & DVECTOR_MASK) && atomKK->k_dvector.need_sync<LMPDeviceType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_2d>(atomKK->k_dvector,space);
+  } else {
+    if ((mask & MOLECULE_MASK) && atomKK->k_molecule.need_sync<LMPHostType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_tagint_1d>(atomKK->k_molecule,space);
+    if ((mask & Q_MASK) && atomKK->k_q.need_sync<LMPHostType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_1d>(atomKK->k_q,space);
+    if ((mask & RMASS_MASK) && atomKK->k_rmass.need_sync<LMPHostType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_1d>(atomKK->k_rmass,space);
+    if ((mask & DVECTOR_MASK) && atomKK->k_dvector.need_sync<LMPHostType>())
+      atomKK->avecKK->perform_async_copy<DAT::tdual_float_2d>(atomKK->k_dvector,space);
+  }
+}
+
+/* ---------------------------------------------------------------------- */
+
+void FixPropertyAtomKokkos::modified(ExecutionSpace space, unsigned int mask)
+{
+  if (space == Device) {
+    if (molecule_flag && (mask & MOLECULE_MASK)) atomKK->k_molecule.modify<LMPDeviceType>();
+    if (q_flag && (mask & Q_MASK)) atomKK->k_q.modify<LMPDeviceType>();
+    if (rmass_flag && (mask & RMASS_MASK)) atomKK->k_rmass.modify<LMPDeviceType>();
+    if (dvector_flag && (mask & DVECTOR_MASK)) atomKK->k_dvector.modify<LMPDeviceType>();
+  } else {
+    if (molecule_flag && (mask & MOLECULE_MASK)) atomKK->k_molecule.modify<LMPHostType>();
+    if (q_flag && (mask & Q_MASK)) atomKK->k_q.modify<LMPHostType>();
+    if (rmass_flag && (mask & RMASS_MASK)) atomKK->k_rmass.modify<LMPHostType>();
+    if (dvector_flag && (mask & DVECTOR_MASK)) atomKK->k_dvector.modify<LMPHostType>();
+  }
+}
--- a/src/KOKKOS/fix_property_atom_kokkos.h
+++ b/src/KOKKOS/fix_property_atom_kokkos.h
@ -22,14 +22,23 @@ FixStyle(property/atom/kk,FixPropertyAtomKokkos);
 #define LMP_FIX_PROPERTY_ATOM_KOKKOS_H

 #include "fix_property_atom.h"
+#include "atom_vec_kokkos.h"

 namespace LAMMPS_NS {

 class FixPropertyAtomKokkos : public FixPropertyAtom {
 public:
  FixPropertyAtomKokkos(class LAMMPS *, int, char **);
-
+  void post_constructor() override;
+  ~FixPropertyAtomKokkos() override;
  void grow_arrays(int) override;
+
+  void sync(ExecutionSpace space, unsigned int mask);
+  void modified(ExecutionSpace space, unsigned int mask);
+  void sync_overlapping_device(ExecutionSpace space, unsigned int mask);
+
+ private:
+  int dvector_flag;
 };

 }
--- a/src/KOKKOS/kokkos.cpp
+++ b/src/KOKKOS/kokkos.cpp
@ -104,7 +104,9 @@ KokkosLMP::KokkosLMP(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)

  int me = 0;
  MPI_Comm_rank(world,&me);
-  if (me == 0) error->message(FLERR,"KOKKOS mode is enabled");
+  if (me == 0)
+    error->message(FLERR,"KOKKOS mode with Kokkos version {}.{}.{} is enabled",
+                   KOKKOS_VERSION / 10000, (KOKKOS_VERSION % 10000) / 100, KOKKOS_VERSION % 100);

  // process any command-line args that invoke Kokkos settings

@ -143,6 +145,14 @@ KokkosLMP::KokkosLMP(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)
        if (device >= skip_gpu) device++;
        set_flag = 1;
      }
+      if ((str = getenv("FLUX_TASK_LOCAL_ID"))) {
+        if (ngpus > 0) {
+          int local_rank = atoi(str);
+          device = local_rank % ngpus;
+          if (device >= skip_gpu) device++;
+          set_flag = 1;
+        }
+      }
      if ((str = getenv("MPT_LRANK"))) {
        if (ngpus > 0) {
          int local_rank = atoi(str);
--- a/src/KOKKOS/kokkos_base.h
+++ b/src/KOKKOS/kokkos_base.h
@ -41,11 +41,6 @@ class KokkosBase {
                                           int, int *) {return 0;};
  virtual void unpack_forward_comm_fix_kokkos(int, int, DAT::tdual_xfloat_1d &) {}

-
-  // Region
-  virtual void match_all_kokkos(int, DAT::tdual_int_1d) {}
-
-  // Fix
  virtual int pack_exchange_kokkos(const int & /*nsend*/, DAT::tdual_xfloat_2d & /*k_buf*/,
                                   DAT::tdual_int_1d /*k_sendlist*/,
                                   DAT::tdual_int_1d /*k_copylist*/,
@ -54,6 +49,9 @@ class KokkosBase {
                                      DAT::tdual_int_1d & /*indices*/, int /*nrecv*/,
                                      ExecutionSpace /*space*/) {}

+  // Region
+  virtual void match_all_kokkos(int, DAT::tdual_int_1d) {}
+
  using KeyViewType = DAT::t_x_array;
  using BinOp = BinOp3DLAMMPS<KeyViewType>;
  virtual void
--- a/src/KOKKOS/min_kokkos.cpp
+++ b/src/KOKKOS/min_kokkos.cpp
@ -59,6 +59,9 @@ void MinKokkos::init()
 {
  Min::init();

+  if (!fix_minimize->kokkosable)
+    error->all(FLERR,"KOKKOS package requires fix minimize/kk");
+
  fix_minimize_kk = (FixMinimizeKokkos*) fix_minimize;
 }

@ -69,8 +72,7 @@ void MinKokkos::init()
 void MinKokkos::setup(int flag)
 {
  if (comm->me == 0 && screen) {
-    fmt::print(screen,"Setting up {} style minimization ...\n",
-               update->minimize_style);
+    fmt::print(screen,"Setting up {} style minimization ...\n", update->minimize_style);
    if (flag) {
      fmt::print(screen,"  Unit style    : {}\n", update->unit_style);
      fmt::print(screen,"  Current step  : {}\n", update->ntimestep);
@ -89,14 +91,13 @@ void MinKokkos::setup(int flag)
    fextra = new double[nextra_global];
    if (comm->me == 0)
      error->warning(FLERR, "Energy due to {} extra global DOFs will"
-                     " be included in minimizer energies\n", nextra_global);
+                     " be included in minimizer energies\n",nextra_global);
  }

  // compute for potential energy

-  int id = modify->find_compute("thermo_pe");
-  if (id < 0) error->all(FLERR,"Minimization could not find thermo_pe compute");
-  pe_compute = modify->compute[id];
+  pe_compute = modify->get_compute_by_id("thermo_pe");
+  if (!pe_compute) error->all(FLERR,"Minimization could not find thermo_pe compute");

  // style-specific setup does two tasks
  // setup extra global dof vectors
@ -534,6 +535,7 @@ double MinKokkos::energy_force(int resetflag)
    if (resetflag) fix_minimize_kk->reset_coords();
    reset_vectors();
  }
+
  return energy;
 }

@ -572,7 +574,14 @@ void MinKokkos::force_clear()
        l_torque(i,2) = 0.0;
      }
    });
+
+    if (extraflag) {
+      size_t nbytes = sizeof(double) * atom->nlocal;
+      if (force->newton) nbytes += sizeof(double) * atom->nghost;
+      atom->avec->force_clear(0,nbytes);
+    }
  }
+
  atomKK->modified(Device,F_MASK|TORQUE_MASK);
 }

--- a/src/KOKKOS/mliap_data_kokkos.cpp
+++ b/src/KOKKOS/mliap_data_kokkos.cpp
@ -85,6 +85,7 @@ void MLIAPDataKokkos<DeviceType>::generate_neighdata(class NeighList *list_in, i
  // clear gradforce and elems arrays

  int nall = atom->nlocal + atom->nghost;
+  nlocal = atom->nlocal;
  ntotal = nall;
  if (gradgradflag > -1){
    auto d_gradforce = k_gradforce.template view<DeviceType>();
--- a/src/KOKKOS/mliap_data_kokkos.h
+++ b/src/KOKKOS/mliap_data_kokkos.h
@ -118,6 +118,7 @@ public:
    egradient(nullptr),
    ntotal(base.ntotal),
    nlistatoms(base.nlistatoms),
+    nlocal(base.nlocal),
    natomneigh(base.natomneigh),
    numneighs(base.numneighs),
    iatoms(base.k_iatoms.d_view.data()),
@ -171,6 +172,7 @@ public:
  // Neighborlist stuff
  const int ntotal;
  const int nlistatoms;
+  const int nlocal;
  const int natomneigh;
  int *numneighs;
  int *iatoms;
@ -191,7 +193,7 @@ public:
  int dev;

 #ifdef LMP_KOKKOS_GPU
-  MLIAPDataKokkosDevice(MLIAPDataKokkos<LMPHostType> &base) : ndescriptors(-1),nparams(-1),nelements(-1),ntotal(-1),nlistatoms(-1),natomneigh(-1),
+  MLIAPDataKokkosDevice(MLIAPDataKokkos<LMPHostType> &base) : ndescriptors(-1),nparams(-1),nelements(-1),ntotal(-1),nlistatoms(-1),nlocal(-1),natomneigh(-1),
      nneigh_max(-1),npairs(-1)
  {
    // It cannot get here, but needed for compilation
--- a/src/KOKKOS/mliap_model_python_couple_kokkos.pyx
+++ b/src/KOKKOS/mliap_model_python_couple_kokkos.pyx
@ -25,6 +25,7 @@ cdef extern from "mliap_data_kokkos.h" namespace "LAMMPS_NS":
    cdef cppclass MLIAPDataKokkosDevice:
        # Array shapes
        int nlistatoms
+        int nlocal
        int ndescriptors

        # Input data
@ -130,14 +131,14 @@ cdef create_array(device, void *pointer, shape,is_int):
                return numpy.asarray(<int[:shape[0],:shape[1]]>pointer)
            else:
                return numpy.asarray(<double[:shape[0],:shape[1]]>pointer)
-                
+

 cdef public void MLIAPPYKokkos_compute_gradients(MLIAPModelPythonKokkosDevice * c_model, MLIAPDataKokkosDevice * data) with gil:

    dev=data.dev

    torch.cuda.nvtx.range_push("set data fields")
-    model = retrieve(c_model) 
+    model = retrieve(c_model)
    n_d = data.ndescriptors
    n_a = data.nlistatoms

@ -148,7 +149,7 @@ cdef public void MLIAPPYKokkos_compute_gradients(MLIAPModelPythonKokkosDevice *
    beta_cp = create_array(dev, data.betas, (n_a, n_d), False)
    desc_cp = create_array(dev, data.descriptors, (n_a, n_d), False)
    torch.cuda.nvtx.range_pop()
-    
+
    # Invoke python model on numpy arrays.
    torch.cuda.nvtx.range_push("call model")
    model(elem_cp,desc_cp,beta_cp,en_cp,dev==1)
--- a/src/KOKKOS/mliap_unified_couple_kokkos.pyx
+++ b/src/KOKKOS/mliap_unified_couple_kokkos.pyx
@ -59,6 +59,7 @@ cdef extern from "mliap_data_kokkos.h" namespace "LAMMPS_NS":

        int ntotal              # total number of owned and ghost atoms on this proc
        int nlistatoms          # current number of atoms in local atom lists
+        int nlocal
        int natomneigh          # current number of atoms and ghosts in atom neighbor arrays
        int * numneighs         # neighbors count for each atom
        int * iatoms            # index of each atom
@ -133,14 +134,14 @@ cdef create_array(device, void *pointer, shape,is_int):
                return np.asarray(<int[:shape[0],:shape[1]]>pointer)
            else:
                return np.asarray(<double[:shape[0],:shape[1]]>pointer)
-            
+


 # Cython implementation of MLIAPData
 # Automatically converts between C arrays and numpy when needed
 cdef class MLIAPDataPy:
    cdef MLIAPDataKokkosDevice * data
-    
+
    def __cinit__(self):
        self.data = NULL

@ -157,7 +158,7 @@ cdef class MLIAPDataPy:
            ptr = eij.data.ptr
        except:
            ptr = eij.data_ptr()
-        update_pair_energy(self.data, <double*>ptr)  
+        update_pair_energy(self.data, <double*>ptr)
    def update_pair_energy(self, eij):
        if self.data.dev==0:
            self.update_pair_energy_cpu(eij)
@ -177,7 +178,7 @@ cdef class MLIAPDataPy:
            ptr = fij.data.ptr
        except:
            ptr = fij.data_ptr()
-        update_pair_forces(self.data, <double*>ptr)  
+        update_pair_forces(self.data, <double*>ptr)
    def update_pair_forces(self, fij):
        if self.data.dev==0:
            self.update_pair_forces_cpu(fij)
@ -189,11 +190,11 @@ cdef class MLIAPDataPy:
            return None
        return create_array(self.data.dev, self.data.f, [self.ntotal, 3],False)

-    
+
    @property
    def size_gradforce(self):
        return self.data.size_gradforce
- 
+
    @write_only_property
    def gradforce(self, value):
        if self.data.gradforce is NULL:
@ -202,7 +203,7 @@ cdef class MLIAPDataPy:
        cdef double[:, :] value_view = value
        gradforce_view[:] = value_view
        print("This code has not been tested or optimized for the GPU, if you are getting this warning optimize gradforce")
- 
+
    @write_only_property
    def betas(self, value):
        if self.data.betas is NULL:
@ -280,7 +281,7 @@ cdef class MLIAPDataPy:
    @property
    def ntotal(self):
        return self.data.ntotal
-    
+
    @property
    def elems(self):
        if self.data.elems is NULL:
@ -290,7 +291,11 @@ cdef class MLIAPDataPy:
    @property
    def nlistatoms(self):
        return self.data.nlistatoms
-    
+
+    @property
+    def nlocal(self):
+        return self.data.nlocal
+
    @property
    def natomneigh(self):
        return self.data.natomneigh
@ -306,7 +311,7 @@ cdef class MLIAPDataPy:
        if self.data.iatoms is NULL:
            return None
        return create_array(self.data.dev, self.data.iatoms, [self.natomneigh],True)
-    
+
    @property
    def ielems(self):
        if self.data.ielems is NULL:
@ -322,7 +327,7 @@ cdef class MLIAPDataPy:
        if self.data.pair_i is NULL:
            return None
        return create_array(self.data.dev, self.data.pair_i, [self.npairs],True)
-    
+
    @property
    def pair_j(self):
        return self.jatoms
@ -332,7 +337,7 @@ cdef class MLIAPDataPy:
        if self.data.jatoms is NULL:
            return None
        return create_array(self.data.dev, self.data.jatoms, [self.npairs],True)
-    
+
    @property
    def jelems(self):
        if self.data.jelems is NULL:
@ -383,13 +388,13 @@ cdef class MLIAPUnifiedInterfaceKokkos:
        self.model = NULL
        self.descriptor = NULL
        self.unified_impl = unified_impl
-    
+
    def compute_gradients(self, data):
        self.unified_impl.compute_gradients(data)
-    
+
    def compute_descriptors(self, data):
        self.unified_impl.compute_descriptors(data)
-    
+
    def compute_forces(self, data):
        self.unified_impl.compute_forces(data)

@ -443,7 +448,7 @@ cdef public object mliap_unified_connect_kokkos(char *fname, MLIAPDummyModel * m

    if unified.element_types is None:
        raise ValueError("no element type set")
-    
+
    cdef int nelements = <int>len(unified.element_types)
    cdef char **elements = <char**>malloc(nelements * sizeof(char*))

--- a/src/KOKKOS/mliap_unified_kokkos.cpp
+++ b/src/KOKKOS/mliap_unified_kokkos.cpp
@ -271,8 +271,8 @@ void LAMMPS_NS::update_pair_energy(MLIAPDataKokkosDevice *data, double *eij)
 {
  auto d_eatoms = data->eatoms;
  auto d_pair_i= data->pair_i;
-  const auto nlistatoms = data->nlistatoms;
-  Kokkos::parallel_for(nlistatoms, KOKKOS_LAMBDA(int ii){
+  const auto nlocal = data->nlocal;
+  Kokkos::parallel_for(nlocal, KOKKOS_LAMBDA(int ii){
    d_eatoms[ii] = 0;
  });

@ -281,7 +281,7 @@ void LAMMPS_NS::update_pair_energy(MLIAPDataKokkosDevice *data, double *eij)
    double e = 0.5 * eij[ii];

    // must not count any contribution where i is not a local atom
-    if (i < nlistatoms) {
+    if (i < nlocal) {
      Kokkos::atomic_add(&d_eatoms[i], e);
      local_sum += e;
    }
@ -294,7 +294,7 @@ void LAMMPS_NS::update_pair_energy(MLIAPDataKokkosDevice *data, double *eij)

 void LAMMPS_NS::update_pair_forces(MLIAPDataKokkosDevice *data, double *fij)
 {
-  const auto nlistatoms = data->nlistatoms;
+  const auto nlocal = data->nlocal;
  auto *f = data->f;
  auto pair_i = data->pair_i;
  auto j_atoms = data->jatoms;
@ -315,7 +315,7 @@ void LAMMPS_NS::update_pair_forces(MLIAPDataKokkosDevice *data, double *fij)
    int i = pair_i[ii];
    int j = j_atoms[ii];
    // must not count any contribution where i is not a local atom
-    if (i < nlistatoms) {
+    if (i < nlocal) {
      Kokkos::atomic_add(&f[i*3+0], fij[ii3+0]);
      Kokkos::atomic_add(&f[i*3+1], fij[ii3+1]);
      Kokkos::atomic_add(&f[i*3+2], fij[ii3+2]);
@ -378,12 +378,12 @@ void LAMMPS_NS::update_pair_forces(MLIAPDataKokkosDevice *data, double *fij)
 void LAMMPS_NS::update_atom_energy(MLIAPDataKokkosDevice *data, double *ei)
 {
  auto d_eatoms = data->eatoms;
-  const auto nlistatoms = data->nlistatoms;
+  const auto nlocal = data->nlocal;

-  Kokkos::parallel_reduce(nlistatoms, KOKKOS_LAMBDA(int i, double &local_sum){
+  Kokkos::parallel_reduce(nlocal, KOKKOS_LAMBDA(int i, double &local_sum){
    double e = ei[i];
    // must not count any contribution where i is not a local atom
-    if (i < nlistatoms) {
+    if (i < nlocal) {
      d_eatoms[i] = e;
      local_sum += e;
    }
--- a/src/KOKKOS/modify_kokkos.cpp
+++ b/src/KOKKOS/modify_kokkos.cpp
@ -362,6 +362,17 @@ void ModifyKokkos::pre_reverse(int eflag, int vflag)

 void ModifyKokkos::post_force(int vflag)
 {
+  for (int i = 0; i < n_post_force_group; i++) {
+    atomKK->sync(fix[list_post_force_group[i]]->execution_space,
+                 fix[list_post_force_group[i]]->datamask_read);
+    int prev_auto_sync = lmp->kokkos->auto_sync;
+    if (!fix[list_post_force_group[i]]->kokkosable) lmp->kokkos->auto_sync = 1;
+    fix[list_post_force_group[i]]->post_force(vflag);
+    lmp->kokkos->auto_sync = prev_auto_sync;
+    atomKK->modified(fix[list_post_force_group[i]]->execution_space,
+                     fix[list_post_force_group[i]]->datamask_modify);
+  }
+
  for (int i = 0; i < n_post_force; i++) {
    atomKK->sync(fix[list_post_force[i]]->execution_space,
                 fix[list_post_force[i]]->datamask_read);
--- a/src/KOKKOS/neigh_bond_kokkos.cpp
+++ b/src/KOKKOS/neigh_bond_kokkos.cpp
@ -112,9 +112,8 @@ void NeighBondKokkos<DeviceType>::init_topology_kk() {
  int i,m;
  int bond_off = 0;
  int angle_off = 0;
-  for (i = 0; i < modify->nfix; i++)
-    if ((strcmp(modify->fix[i]->style,"shake") == 0)
-        || (strcmp(modify->fix[i]->style,"rattle") == 0))
+  for (const auto &ifix : modify->get_fix_list())
+    if (utils::strmatch(ifix->style,"^shake") || utils::strmatch(ifix->style,"^rattle"))
      bond_off = angle_off = 1;
  if (force->bond && force->bond_match("quartic")) bond_off = 1;

--- a/src/KOKKOS/neighbor_kokkos.cpp
+++ b/src/KOKKOS/neighbor_kokkos.cpp
@ -308,7 +308,8 @@ void NeighborKokkos::build_kokkos(int topoflag)
  for (i = 0; i < npair_perpetual; i++) {
    m = plist[i];
    if (!lists[m]->kokkos) atomKK->sync(Host,ALL_MASK);
-    if (!lists[m]->copy) lists[m]->grow(nlocal,nall);
+    if (!lists[m]->copy || lists[m]->trim || lists[m]->kk2cpu)
+      lists[m]->grow(nlocal,nall);
    neigh_pair[m]->build_setup();
    neigh_pair[m]->build(lists[m]);
  }
--- a/src/KOKKOS/npair_kokkos.cpp
+++ b/src/KOKKOS/npair_kokkos.cpp
@ -113,7 +113,7 @@ void NPairKokkos<DeviceType,HALF,NEWTON,GHOST,TRI,SIZE>::copy_stencil_info()
  NPair::copy_stencil_info();
  nstencil = ns->nstencil;

-  if (ns->last_stencil != last_stencil_old) {
+  if (ns->last_stencil != last_stencil_old || ns->last_stencil == update->ntimestep) {
    // copy stencil to device as it may have changed

    last_stencil_old = ns->last_stencil;
--- a/src/KOKKOS/npair_trim_kokkos.cpp
+++ b/src/KOKKOS/npair_trim_kokkos.cpp
@ -62,8 +62,8 @@ void NPairTrimKokkos<DeviceType>::trim_to_kokkos(NeighList *list)
  d_ilist_copy = k_list_copy->d_ilist;
  d_numneigh_copy = k_list_copy->d_numneigh;
  d_neighbors_copy = k_list_copy->d_neighbors;
-  int inum_copy = list->listcopy->inum;
-  if (list->ghost) inum_copy += list->listcopy->gnum;
+  int inum_trim = list->listcopy->inum;
+  if (list->ghost) inum_trim += list->listcopy->gnum;

  NeighListKokkos<DeviceType>* k_list = static_cast<NeighListKokkos<DeviceType>*>(list);
  k_list->maxneighs = k_list_copy->maxneighs; // simple, but could be made more memory efficient
@ -75,7 +75,7 @@ void NPairTrimKokkos<DeviceType>::trim_to_kokkos(NeighList *list)
  // loop over parent list and trim

  copymode = 1;
-  Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagNPairTrim>(0,inum_copy),*this);
+  Kokkos::parallel_for(Kokkos::RangePolicy<DeviceType, TagNPairTrim>(0,inum_trim),*this);
  copymode = 0;

  list->inum = k_list_copy->inum;
@ -132,8 +132,8 @@ void NPairTrimKokkos<DeviceType>::trim_to_cpu(NeighList *list)

  int inum = listcopy->inum;
  int gnum = listcopy->gnum;
-  int inum_all = inum;
-  if (list->ghost) inum_all += gnum;
+  int inum_trim = inum;
+  if (list->ghost) inum_trim += gnum;
  auto h_ilist = listcopy_kk->k_ilist.h_view;
  auto h_numneigh = Kokkos::create_mirror_view_and_copy(LMPHostType(),listcopy_kk->d_numneigh);
  auto h_neighbors = Kokkos::create_mirror_view_and_copy(LMPHostType(),listcopy_kk->d_neighbors);
@ -151,7 +151,7 @@ void NPairTrimKokkos<DeviceType>::trim_to_cpu(NeighList *list)
  MyPage<int> *ipage = list->ipage;
  ipage->reset();

-  for (int ii = 0; ii < inum_all; ii++) {
+  for (int ii = 0; ii < inum_trim; ii++) {
    int n = 0;
    neighptr = ipage->vget();

--- a/src/KOKKOS/pair_buck_coul_cut_kokkos.h
+++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.h
@ -112,15 +112,18 @@ class PairBuckCoulCutKokkos : public PairBuckCoulCut {

  void allocate() override;

-  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairBuckCoulCutKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,FULL,void>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALF,void>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALFTHREAD,void>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,FULL,0>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,FULL,1>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALF>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulCutKokkos,HALFTHREAD>(PairBuckCoulCutKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairBuckCoulCutKokkos,void>(PairBuckCoulCutKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairBuckCoulCutKokkos>(PairBuckCoulCutKokkos*);
--- a/src/KOKKOS/pair_buck_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_buck_coul_long_kokkos.h
@ -115,27 +115,33 @@ class PairBuckCoulLongKokkos : public PairBuckCoulLong {

  void allocate() override;

-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,CoulLongTable<1> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALF,CoulLongTable<1> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairBuckCoulLongKokkos,CoulLongTable<1> >(PairBuckCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,0,CoulLongTable<1>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,1,CoulLongTable<1>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALF,0,CoulLongTable<1>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairBuckCoulLongKokkos,CoulLongTable<1>>(PairBuckCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,CoulLongTable<0> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALF,CoulLongTable<0> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairBuckCoulLongKokkos,CoulLongTable<0> >(PairBuckCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairBuckCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,0,CoulLongTable<0>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,FULL,1,CoulLongTable<0>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALF,0,CoulLongTable<0>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckCoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairBuckCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairBuckCoulLongKokkos,CoulLongTable<0>>(PairBuckCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairBuckCoulLongKokkos>(PairBuckCoulLongKokkos*);

--- a/src/KOKKOS/pair_buck_kokkos.h
+++ b/src/KOKKOS/pair_buck_kokkos.h
@ -91,16 +91,19 @@ class PairBuckKokkos : public PairBuck {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairBuckKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairBuckKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairBuckKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairBuckKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairBuckKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairBuckKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairBuckKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairBuckKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairBuckKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairBuckKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,FULL,void>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,HALF,void>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,HALFTHREAD,void>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairBuckKokkos,void>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,FULL,0>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,FULL,1>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,HALF>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairBuckKokkos,HALFTHREAD>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairBuckKokkos>(PairBuckKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairBuckKokkos>(PairBuckKokkos*);
 };

--- a/src/KOKKOS/pair_coul_cut_kokkos.h
+++ b/src/KOKKOS/pair_coul_cut_kokkos.h
@ -112,15 +112,18 @@ class PairCoulCutKokkos : public PairCoulCut {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairCoulCutKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairCoulCutKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairCoulCutKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairCoulCutKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairCoulCutKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,FULL,void>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,HALF,void>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,HALFTHREAD,void>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,FULL,0>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,FULL,1>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,HALF>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulCutKokkos,HALFTHREAD>(PairCoulCutKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairCoulCutKokkos,void>(PairCoulCutKokkos*,
                                                       NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairCoulCutKokkos>(PairCoulCutKokkos*);
--- a/src/KOKKOS/pair_coul_debye_kokkos.h
+++ b/src/KOKKOS/pair_coul_debye_kokkos.h
@ -112,15 +112,18 @@ class PairCoulDebyeKokkos : public PairCoulDebye {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairCoulDebyeKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairCoulDebyeKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairCoulDebyeKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairCoulDebyeKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairCoulDebyeKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,FULL,void>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALF,void>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALFTHREAD,void>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,FULL,0>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,FULL,1>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALF>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulDebyeKokkos,HALFTHREAD>(PairCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairCoulDebyeKokkos,void>(PairCoulDebyeKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairCoulDebyeKokkos>(PairCoulDebyeKokkos*);
--- a/src/KOKKOS/pair_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_coul_long_kokkos.h
@ -114,27 +114,33 @@ class PairCoulLongKokkos : public PairCoulLong {

  void allocate() override;

-  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,CoulLongTable<1> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALF,CoulLongTable<1> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairCoulLongKokkos,CoulLongTable<1> >(PairCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,0,CoulLongTable<1>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,1,CoulLongTable<1>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALF,0,CoulLongTable<1>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairCoulLongKokkos,CoulLongTable<1>>(PairCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,CoulLongTable<0> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALF,CoulLongTable<0> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairCoulLongKokkos,CoulLongTable<0> >(PairCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,0,CoulLongTable<0>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,FULL,1,CoulLongTable<0>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALF,0,CoulLongTable<0>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairCoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairCoulLongKokkos,CoulLongTable<0>>(PairCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairCoulLongKokkos>(PairCoulLongKokkos*);

--- a/src/KOKKOS/pair_eam_alloy_kokkos.cpp
+++ b/src/KOKKOS/pair_eam_alloy_kokkos.cpp
@ -1477,7 +1477,7 @@ void PairEAMAlloyKokkos<DeviceType>::file2array_alloy()
 template<typename DeviceType>
 template<class TAG>
 struct PairEAMAlloyKokkos<DeviceType>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    auto policy = Kokkos::RangePolicy<DeviceType, TAG>(0,inum);
    return policy;
@ -1488,7 +1488,7 @@ struct PairEAMAlloyKokkos<DeviceType>::policyInstance {
 template<>
 template<class TAG>
 struct PairEAMAlloyKokkos<Kokkos::Experimental::HIP>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    static_assert(t_ffloat_2d_n7::static_extent(2) == 7,
                  "Breaking assumption of spline dim for KernelAB and KernelC scratch caching");
--- a/src/KOKKOS/pair_eam_fs_kokkos.cpp
+++ b/src/KOKKOS/pair_eam_fs_kokkos.cpp
@ -1487,7 +1487,7 @@ void PairEAMFSKokkos<DeviceType>::file2array_fs()
 template<typename DeviceType>
 template<class TAG>
 struct PairEAMFSKokkos<DeviceType>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    auto policy = Kokkos::RangePolicy<DeviceType, TAG>(0,inum);
    return policy;
@ -1498,7 +1498,7 @@ struct PairEAMFSKokkos<DeviceType>::policyInstance {
 template<>
 template<class TAG>
 struct PairEAMFSKokkos<Kokkos::Experimental::HIP>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    static_assert(t_ffloat_2d_n7::static_extent(2) == 7,
                  "Breaking assumption of spline dim for KernelAB and KernelC scratch caching");
--- a/src/KOKKOS/pair_eam_kokkos.cpp
+++ b/src/KOKKOS/pair_eam_kokkos.cpp
@ -1162,7 +1162,7 @@ void PairEAMKokkos<DeviceType>::ev_tally(EV_FLOAT &ev, const int &i, const int &
 template<typename DeviceType>
 template<class TAG>
 struct PairEAMKokkos<DeviceType>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    auto policy = Kokkos::RangePolicy<DeviceType, TAG>(0,inum);
    return policy;
@ -1173,7 +1173,7 @@ struct PairEAMKokkos<DeviceType>::policyInstance {
 template<>
 template<class TAG>
 struct PairEAMKokkos<Kokkos::Experimental::HIP>::policyInstance {
-  KOKKOS_INLINE_FUNCTION
+
  static auto get(int inum) {
    static_assert(t_ffloat_2d_n7::static_extent(2) == 7,
                  "Breaking assumption of spline dim for KernelAB and KernelC scratch caching");
--- a/src/KOKKOS/pair_kokkos.h
+++ b/src/KOKKOS/pair_kokkos.h
@ -50,7 +50,7 @@ struct DoCoul<1> {


 //Specialisation for Neighborlist types Half, HalfThread, Full
-template <class PairStyle, int NEIGHFLAG, bool STACKPARAMS, class Specialisation = void>
+template <class PairStyle, int NEIGHFLAG, bool STACKPARAMS, int ZEROFLAG = 0, class Specialisation = void>
 struct PairComputeFunctor  {
  typedef typename PairStyle::device_type device_type ;
  typedef ArrayTypes<device_type> AT;
@ -137,7 +137,7 @@ struct PairComputeFunctor  {
    F_FLOAT fytmp = 0.0;
    F_FLOAT fztmp = 0.0;

-    if (NEIGHFLAG == FULL) {
+    if (NEIGHFLAG == FULL && ZEROFLAG) {
      f(i,0) = 0.0;
      f(i,1) = 0.0;
      f(i,2) = 0.0;
@ -211,7 +211,7 @@ struct PairComputeFunctor  {
    F_FLOAT fytmp = 0.0;
    F_FLOAT fztmp = 0.0;

-    if (NEIGHFLAG == FULL) {
+    if (NEIGHFLAG == FULL && ZEROFLAG) {
      f(i,0) = 0.0;
      f(i,1) = 0.0;
      f(i,2) = 0.0;
@ -292,11 +292,13 @@ struct PairComputeFunctor  {
      const X_FLOAT ztmp = c.x(i,2);
      const int itype = c.type(i);

-      Kokkos::single(Kokkos::PerThread(team), [&] (){
-        f(i,0) = 0.0;
-        f(i,1) = 0.0;
-        f(i,2) = 0.0;
-      });
+      if (ZEROFLAG) {
+        Kokkos::single(Kokkos::PerThread(team), [&] (){
+          f(i,0) = 0.0;
+          f(i,1) = 0.0;
+          f(i,2) = 0.0;
+        });
+      }

      const AtomNeighborsConst neighbors_i = list.get_neighbors_const(i);
      const int jnum = list.d_numneigh[i];
@ -355,11 +357,13 @@ struct PairComputeFunctor  {
      const int itype = c.type(i);
      const F_FLOAT qtmp = c.q(i);

-      Kokkos::single(Kokkos::PerThread(team), [&] (){
-        f(i,0) = 0.0;
-        f(i,1) = 0.0;
-        f(i,2) = 0.0;
-      });
+      if (ZEROFLAG) {
+        Kokkos::single(Kokkos::PerThread(team), [&] (){
+          f(i,0) = 0.0;
+          f(i,1) = 0.0;
+          f(i,2) = 0.0;
+        });
+      }

      const AtomNeighborsConst neighbors_i = list.get_neighbors_const(i);
      const int jnum = list.d_numneigh[i];
@ -423,11 +427,13 @@ struct PairComputeFunctor  {
      const X_FLOAT ztmp = c.x(i,2);
      const int itype = c.type(i);

-      Kokkos::single(Kokkos::PerThread(team), [&] (){
-        f(i,0) = 0.0;
-        f(i,1) = 0.0;
-        f(i,2) = 0.0;
-      });
+      if (ZEROFLAG) {
+        Kokkos::single(Kokkos::PerThread(team), [&] (){
+          f(i,0) = 0.0;
+          f(i,1) = 0.0;
+          f(i,2) = 0.0;
+        });
+      }

      const AtomNeighborsConst neighbors_i = list.get_neighbors_const(i);
      const int jnum = list.d_numneigh[i];
@ -525,11 +531,13 @@ struct PairComputeFunctor  {
      const int itype = c.type(i);
      const F_FLOAT qtmp = c.q(i);

-      Kokkos::single(Kokkos::PerThread(team), [&] (){
-        f(i,0) = 0.0;
-        f(i,1) = 0.0;
-        f(i,2) = 0.0;
-      });
+      if (ZEROFLAG) {
+        Kokkos::single(Kokkos::PerThread(team), [&] (){
+          f(i,0) = 0.0;
+          f(i,1) = 0.0;
+          f(i,2) = 0.0;
+        });
+      }

      const AtomNeighborsConst neighbors_i = list.get_neighbors_const(i);
      const int jnum = list.d_numneigh[i];
@ -740,7 +748,7 @@ struct PairComputeFunctor  {
 // By having the enable_if with a ! and without it, exactly one of the functions
 // pair_compute_neighlist will match - either the dummy version
 // or the real one further below.
-template<class PairStyle, unsigned NEIGHFLAG, class Specialisation>
+template<class PairStyle, unsigned NEIGHFLAG, int ZEROFLAG=0, class Specialisation = void>
 EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename std::enable_if<!((NEIGHFLAG&PairStyle::EnabledNeighFlags) != 0), NeighListKokkos<typename PairStyle::device_type>*>::type list) {
  EV_FLOAT ev;
  (void) fpair;
@ -770,7 +778,7 @@ int GetTeamSize(FunctorStyle& KOKKOS_GPU_ARG(functor), int KOKKOS_GPU_ARG(inum),
 }

 // Submit ParallelFor for NEIGHFLAG=HALF,HALFTHREAD,FULL
-template<class PairStyle, unsigned NEIGHFLAG, class Specialisation>
+template<class PairStyle, unsigned NEIGHFLAG, int ZEROFLAG = 0, class Specialisation = void>
 EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename std::enable_if<(NEIGHFLAG&PairStyle::EnabledNeighFlags) != 0, NeighListKokkos<typename PairStyle::device_type>*>::type list) {
  EV_FLOAT ev;

@ -784,13 +792,13 @@ EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename std::enable_if<(NEIG
    int atoms_per_team = 32;

    if (fpair->atom->ntypes > MAX_TYPES_STACKPARAMS) {
-      PairComputeFunctor<PairStyle,NEIGHFLAG,false,Specialisation > ff(fpair,list);
+      PairComputeFunctor<PairStyle,NEIGHFLAG,false,ZEROFLAG,Specialisation > ff(fpair,list);
      atoms_per_team = GetTeamSize<typename PairStyle::device_type>(ff, list->inum, (fpair->eflag || fpair->vflag), atoms_per_team, vector_length);
      Kokkos::TeamPolicy<typename PairStyle::device_type,Kokkos::IndexType<int> > policy(list->inum,atoms_per_team,vector_length);
      if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(policy,ff,ev);
      else                              Kokkos::parallel_for(policy,ff);
    } else {
-      PairComputeFunctor<PairStyle,NEIGHFLAG,true,Specialisation > ff(fpair,list);
+      PairComputeFunctor<PairStyle,NEIGHFLAG,true,ZEROFLAG,Specialisation > ff(fpair,list);
      atoms_per_team = GetTeamSize<typename PairStyle::device_type>(ff, list->inum, (fpair->eflag || fpair->vflag), atoms_per_team, vector_length);
      Kokkos::TeamPolicy<typename PairStyle::device_type,Kokkos::IndexType<int> > policy(list->inum,atoms_per_team,vector_length);
      if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(policy,ff,ev);
@ -798,12 +806,12 @@ EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename std::enable_if<(NEIG
    }
  } else {
    if (fpair->atom->ntypes > MAX_TYPES_STACKPARAMS) {
-      PairComputeFunctor<PairStyle,NEIGHFLAG,false,Specialisation > ff(fpair,list);
+      PairComputeFunctor<PairStyle,NEIGHFLAG,false,ZEROFLAG,Specialisation > ff(fpair,list);
      if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(list->inum,ff,ev);
      else                              Kokkos::parallel_for(list->inum,ff);
      ff.contribute();
    } else {
-      PairComputeFunctor<PairStyle,NEIGHFLAG,true,Specialisation > ff(fpair,list);
+      PairComputeFunctor<PairStyle,NEIGHFLAG,true,ZEROFLAG,Specialisation > ff(fpair,list);
      if (fpair->eflag || fpair->vflag) Kokkos::parallel_reduce(list->inum,ff,ev);
      else                              Kokkos::parallel_for(list->inum,ff);
      ff.contribute();
@ -812,16 +820,21 @@ EV_FLOAT pair_compute_neighlist (PairStyle* fpair, typename std::enable_if<(NEIG
  return ev;
 }

-template<class PairStyle, class Specialisation>
+template<class PairStyle, class Specialisation = void>
 EV_FLOAT pair_compute (PairStyle* fpair, NeighListKokkos<typename PairStyle::device_type>* list) {
  EV_FLOAT ev;
  if (fpair->neighflag == FULL) {
-    fpair->fuse_force_clear_flag = 1;
-    ev = pair_compute_neighlist<PairStyle,FULL,Specialisation> (fpair,list);
+    if (utils::strmatch(fpair->lmp->force->pair_style,"^hybrid/overlay")) {
+      fpair->fuse_force_clear_flag = 0;
+      ev = pair_compute_neighlist<PairStyle,FULL,0,Specialisation> (fpair,list);
+    } else {
+      fpair->fuse_force_clear_flag = 1;
+      ev = pair_compute_neighlist<PairStyle,FULL,1,Specialisation> (fpair,list);
+    }
  } else if (fpair->neighflag == HALFTHREAD) {
-    ev = pair_compute_neighlist<PairStyle,HALFTHREAD,Specialisation> (fpair,list);
+    ev = pair_compute_neighlist<PairStyle,HALFTHREAD,0,Specialisation> (fpair,list);
  } else if (fpair->neighflag == HALF) {
-    ev = pair_compute_neighlist<PairStyle,HALF,Specialisation> (fpair,list);
+    ev = pair_compute_neighlist<PairStyle,HALF,0,Specialisation> (fpair,list);
  }
  return ev;
 }
--- a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h
+++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h
@ -110,27 +110,33 @@ class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplic

  void allocate() override;

-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,CoulLongTable<1> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALF,CoulLongTable<1> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmImplicitKokkos,CoulLongTable<1> >(PairLJCharmmCoulCharmmImplicitKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,1,CoulLongTable<1>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALF,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmImplicitKokkos,CoulLongTable<1>>(PairLJCharmmCoulCharmmImplicitKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,CoulLongTable<0> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALF,CoulLongTable<0> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmImplicitKokkos,CoulLongTable<0> >(PairLJCharmmCoulCharmmImplicitKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,FULL,1,CoulLongTable<0>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALF,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmImplicitKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmImplicitKokkos,CoulLongTable<0>>(PairLJCharmmCoulCharmmImplicitKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCharmmCoulCharmmImplicitKokkos>(PairLJCharmmCoulCharmmImplicitKokkos*);

--- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h
+++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h
@ -108,27 +108,33 @@ class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm {

  void allocate() override;

-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,CoulLongTable<1> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALF,CoulLongTable<1> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmKokkos,CoulLongTable<1> >(PairLJCharmmCoulCharmmKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,1,CoulLongTable<1>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALF,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmKokkos,CoulLongTable<1>>(PairLJCharmmCoulCharmmKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,CoulLongTable<0> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALF,CoulLongTable<0> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmKokkos,CoulLongTable<0> >(PairLJCharmmCoulCharmmKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,FULL,1,CoulLongTable<0>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALF,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulCharmmKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulCharmmKokkos,CoulLongTable<0>>(PairLJCharmmCoulCharmmKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCharmmCoulCharmmKokkos>(PairLJCharmmCoulCharmmKokkos*);

--- a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h
@ -106,27 +106,33 @@ class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong {

  void allocate() override;

-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,CoulLongTable<1> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALF,CoulLongTable<1> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulLongKokkos,CoulLongTable<1> >(PairLJCharmmCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,0,CoulLongTable<1>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,1,CoulLongTable<1>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALF,0,CoulLongTable<1>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulLongKokkos,CoulLongTable<1>>(PairLJCharmmCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,CoulLongTable<0> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALF,CoulLongTable<0> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCharmmCoulLongKokkos,CoulLongTable<0> >(PairLJCharmmCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCharmmCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,0,CoulLongTable<0>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,FULL,1,CoulLongTable<0>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALF,0,CoulLongTable<0>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCharmmCoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJCharmmCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCharmmCoulLongKokkos,CoulLongTable<0>>(PairLJCharmmCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCharmmCoulLongKokkos>(PairLJCharmmCoulLongKokkos*);

--- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h
+++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h
@ -104,15 +104,18 @@ class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJClass2CoulCutKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,FULL,void>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALF,void>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALFTHREAD,void>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,FULL,0>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,FULL,1>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALF>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulCutKokkos,HALFTHREAD>(PairLJClass2CoulCutKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairLJClass2CoulCutKokkos,void>(PairLJClass2CoulCutKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJClass2CoulCutKokkos>(PairLJClass2CoulCutKokkos*);
--- a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h
@ -107,27 +107,33 @@ class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,CoulLongTable<1> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALF,CoulLongTable<1> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJClass2CoulLongKokkos,CoulLongTable<1> >(PairLJClass2CoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,0,CoulLongTable<1>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,1,CoulLongTable<1>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALF,0,CoulLongTable<1>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJClass2CoulLongKokkos,CoulLongTable<1>>(PairLJClass2CoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,CoulLongTable<0> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALF,CoulLongTable<0> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJClass2CoulLongKokkos,CoulLongTable<0> >(PairLJClass2CoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJClass2CoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,0,CoulLongTable<0>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,FULL,1,CoulLongTable<0>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALF,0,CoulLongTable<0>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2CoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJClass2CoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJClass2CoulLongKokkos,CoulLongTable<0>>(PairLJClass2CoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJClass2CoulLongKokkos>(PairLJClass2CoulLongKokkos*);

--- a/src/KOKKOS/pair_lj_class2_kokkos.h
+++ b/src/KOKKOS/pair_lj_class2_kokkos.h
@ -96,16 +96,19 @@ class PairLJClass2Kokkos : public PairLJClass2 {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJClass2Kokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJClass2Kokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJClass2Kokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJClass2Kokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJClass2Kokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,FULL,void>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,HALF,void>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,HALFTHREAD,void>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJClass2Kokkos,void>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,FULL,0>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,FULL,1>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,HALF>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJClass2Kokkos,HALFTHREAD>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJClass2Kokkos>(PairLJClass2Kokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJClass2Kokkos>(PairLJClass2Kokkos*);
 };

--- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h
+++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h
@ -104,15 +104,18 @@ class PairLJCutCoulCutKokkos : public PairLJCutCoulCut {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJCutCoulCutKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,FULL,void>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALF,void>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALFTHREAD,void>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,FULL,0>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,FULL,1>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALF>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulCutKokkos,HALFTHREAD>(PairLJCutCoulCutKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairLJCutCoulCutKokkos,void>(PairLJCutCoulCutKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCutCoulCutKokkos>(PairLJCutCoulCutKokkos*);
--- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h
+++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h
@ -104,15 +104,18 @@ class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJCutCoulDebyeKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,FULL,void>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALF,void>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALFTHREAD,void>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,FULL,0>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,FULL,1>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALF>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDebyeKokkos,HALFTHREAD>(PairLJCutCoulDebyeKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairLJCutCoulDebyeKokkos,void>(PairLJCutCoulDebyeKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCutCoulDebyeKokkos>(PairLJCutCoulDebyeKokkos*);
--- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h
+++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h
@ -101,15 +101,18 @@ class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJCutCoulDSFKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,FULL,void>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALF,void>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALFTHREAD,void>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,FULL,0>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,FULL,1>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALF>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulDSFKokkos,HALFTHREAD>(PairLJCutCoulDSFKokkos*,NeighListKokkos<DeviceType>*);
  friend EV_FLOAT pair_compute<PairLJCutCoulDSFKokkos,void>(PairLJCutCoulDSFKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCutCoulDSFKokkos>(PairLJCutCoulDSFKokkos*);
--- a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h
@ -107,27 +107,33 @@ class PairLJCutCoulLongKokkos : public PairLJCutCoulLong {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,CoulLongTable<1> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALF,CoulLongTable<1> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCutCoulLongKokkos,CoulLongTable<1> >(PairLJCutCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,0,CoulLongTable<1>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,1,CoulLongTable<1>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALF,0,CoulLongTable<1>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCutCoulLongKokkos,CoulLongTable<1>>(PairLJCutCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,CoulLongTable<0> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALF,CoulLongTable<0> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCutCoulLongKokkos,CoulLongTable<0> >(PairLJCutCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJCutCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,0,CoulLongTable<0>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,FULL,1,CoulLongTable<0>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALF,0,CoulLongTable<0>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutCoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJCutCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCutCoulLongKokkos,CoulLongTable<0>>(PairLJCutCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCutCoulLongKokkos>(PairLJCutCoulLongKokkos*);

--- a/src/KOKKOS/pair_lj_cut_kokkos.h
+++ b/src/KOKKOS/pair_lj_cut_kokkos.h
@ -92,16 +92,19 @@ class PairLJCutKokkos : public PairLJCut {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJCutKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJCutKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJCutKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJCutKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJCutKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,FULL,void>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,HALF,void>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,HALFTHREAD,void>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJCutKokkos,void>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,FULL,0>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,FULL,1>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,HALF>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJCutKokkos,HALFTHREAD>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJCutKokkos>(PairLJCutKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJCutKokkos>(PairLJCutKokkos*);
 };

--- a/src/KOKKOS/pair_lj_expand_coul_long_kokkos.h
+++ b/src/KOKKOS/pair_lj_expand_coul_long_kokkos.h
@ -116,27 +116,33 @@ class PairLJExpandCoulLongKokkos : public PairLJExpandCoulLong {
  double qqrd2e;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,CoulLongTable<1> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALF,CoulLongTable<1> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJExpandCoulLongKokkos,CoulLongTable<1> >(PairLJExpandCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,0,CoulLongTable<1>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,1,CoulLongTable<1>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALF,0,CoulLongTable<1>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJExpandCoulLongKokkos,CoulLongTable<1>>(PairLJExpandCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,CoulLongTable<0> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALF,CoulLongTable<0> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJExpandCoulLongKokkos,CoulLongTable<0> >(PairLJExpandCoulLongKokkos*,
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJExpandCoulLongKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,0,CoulLongTable<0>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,FULL,1,CoulLongTable<0>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALF,0,CoulLongTable<0>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandCoulLongKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJExpandCoulLongKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJExpandCoulLongKokkos,CoulLongTable<0>>(PairLJExpandCoulLongKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJExpandCoulLongKokkos>(PairLJExpandCoulLongKokkos*);
 };
--- a/src/KOKKOS/pair_lj_expand_kokkos.h
+++ b/src/KOKKOS/pair_lj_expand_kokkos.h
@ -97,16 +97,19 @@ class PairLJExpandKokkos : public PairLJExpand {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJExpandKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJExpandKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJExpandKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJExpandKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJExpandKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,FULL,void>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,HALF,void>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,HALFTHREAD,void>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJExpandKokkos,void>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,FULL,0>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,FULL,1>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,HALF>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJExpandKokkos,HALFTHREAD>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJExpandKokkos>(PairLJExpandKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJExpandKokkos>(PairLJExpandKokkos*);
 };

--- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h
+++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h
@ -115,27 +115,33 @@ class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs {

  void allocate() override;

-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,CoulLongTable<1> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALF,CoulLongTable<1> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJGromacsCoulGromacsKokkos,CoulLongTable<1> >(PairLJGromacsCoulGromacsKokkos*,
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,0,CoulLongTable<1>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,1,CoulLongTable<1>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALF,0,CoulLongTable<1>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJGromacsCoulGromacsKokkos,CoulLongTable<1>>(PairLJGromacsCoulGromacsKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,CoulLongTable<0> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALF,CoulLongTable<0> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJGromacsCoulGromacsKokkos,CoulLongTable<0> >(PairLJGromacsCoulGromacsKokkos*,
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,0,CoulLongTable<0>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,FULL,1,CoulLongTable<0>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALF,0,CoulLongTable<0>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsCoulGromacsKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJGromacsCoulGromacsKokkos,CoulLongTable<0>>(PairLJGromacsCoulGromacsKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJGromacsCoulGromacsKokkos>(PairLJGromacsCoulGromacsKokkos*);

--- a/src/KOKKOS/pair_lj_gromacs_kokkos.h
+++ b/src/KOKKOS/pair_lj_gromacs_kokkos.h
@ -115,27 +115,33 @@ class PairLJGromacsKokkos : public PairLJGromacs {

  void allocate() override;

-  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,true,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,false,CoulLongTable<1> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,false,CoulLongTable<1> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,CoulLongTable<1> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALF,CoulLongTable<1> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALFTHREAD,CoulLongTable<1> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJGromacsKokkos,CoulLongTable<1> >(PairLJGromacsKokkos*,
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,true,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,1,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,false,0,CoulLongTable<1>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,false,0,CoulLongTable<1>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,0,CoulLongTable<1>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,1,CoulLongTable<1>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALF,0,CoulLongTable<1>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALFTHREAD,0,CoulLongTable<1>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJGromacsKokkos,CoulLongTable<1>>(PairLJGromacsKokkos*,
                                                            NeighListKokkos<DeviceType>*);
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,true,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,false,CoulLongTable<0> >;
-  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,false,CoulLongTable<0> >;
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,CoulLongTable<0> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALF,CoulLongTable<0> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALFTHREAD,CoulLongTable<0> >(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJGromacsKokkos,CoulLongTable<0> >(PairLJGromacsKokkos*,
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,true,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,true,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,FULL,false,1,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALF,false,0,CoulLongTable<0>>;
+  friend struct PairComputeFunctor<PairLJGromacsKokkos,HALFTHREAD,false,0,CoulLongTable<0>>;
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,0,CoulLongTable<0>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,FULL,1,CoulLongTable<0>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALF,0,CoulLongTable<0>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJGromacsKokkos,HALFTHREAD,0,CoulLongTable<0>>(PairLJGromacsKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJGromacsKokkos,CoulLongTable<0>>(PairLJGromacsKokkos*,
                                                            NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJGromacsKokkos>(PairLJGromacsKokkos*);

--- a/src/KOKKOS/pair_lj_spica_kokkos.h
+++ b/src/KOKKOS/pair_lj_spica_kokkos.h
@ -97,16 +97,19 @@ class PairLJSPICAKokkos : public PairLJSPICA {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairLJSPICAKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairLJSPICAKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairLJSPICAKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairLJSPICAKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairLJSPICAKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,FULL,void>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,HALF,void>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,HALFTHREAD,void>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairLJSPICAKokkos,void>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,FULL,0>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,FULL,1>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,HALF>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairLJSPICAKokkos,HALFTHREAD>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairLJSPICAKokkos>(PairLJSPICAKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairLJSPICAKokkos>(PairLJSPICAKokkos*);
 };

--- a/src/KOKKOS/pair_mliap_kokkos.cpp
+++ b/src/KOKKOS/pair_mliap_kokkos.cpp
@ -138,6 +138,7 @@ template<class DeviceType>
 void PairMLIAPKokkos<DeviceType>::allocate()
 {
  int n = atom->ntypes;
+
  memoryKK->destroy_kokkos(k_map, map);
  memoryKK->destroy_kokkos(k_cutsq, cutsq);
  memoryKK->destroy_kokkos(k_setflag, setflag);
@ -275,7 +276,10 @@ void PairMLIAPKokkos<DeviceType>::coeff(int narg, char **arg) {
  auto h_cutsq=k_cutsq.template view<LMPHostType>();
  for (int itype=1; itype <= atom->ntypes; ++itype)
    for (int jtype=1; jtype <= atom->ntypes; ++jtype)
-      h_cutsq(itype,jtype) = descriptor->cutsq[map[itype]][map[jtype]];
+      // do not set cuts for NULL atoms
+      if (map[itype] >= 0 && map[jtype] >= 0) {
+         h_cutsq(itype,jtype) = descriptor->cutsq[map[itype]][map[jtype]];
+       }
  k_cutsq.modify<LMPHostType>();
  k_cutsq.sync<DeviceType>();
  constexpr int gradgradflag = -1;
--- a/src/KOKKOS/pair_morse_kokkos.h
+++ b/src/KOKKOS/pair_morse_kokkos.h
@ -92,16 +92,19 @@ class PairMorseKokkos : public PairMorse {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairMorseKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairMorseKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairMorseKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairMorseKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairMorseKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairMorseKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairMorseKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairMorseKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairMorseKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairMorseKokkos,HALFTHREAD,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<PairMorseKokkos,void>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,FULL,0>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,FULL,1>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALF>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairMorseKokkos,HALFTHREAD>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairMorseKokkos>(PairMorseKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairMorseKokkos>(PairMorseKokkos*);
 };

--- a/src/KOKKOS/pair_pace_kokkos.cpp
+++ b/src/KOKKOS/pair_pace_kokkos.cpp
@ -237,6 +237,9 @@ void PairPACEKokkos<DeviceType>::copy_splines()

  ACERadialFunctions* radial_functions = dynamic_cast<ACERadialFunctions*>(basis_set->radial_functions);

+  if (radial_functions == nullptr)
+    error->all(FLERR,"Chosen radial basis style not supported by pair style pace/kk");
+
  for (int i = 0; i < nelements; i++) {
    for (int j = 0; j < nelements; j++) {
      k_splines_gk.h_view(i, j) = radial_functions->splines_gk(i, j);
--- a/src/KOKKOS/pair_snap_kokkos_impl.h
+++ b/src/KOKKOS/pair_snap_kokkos_impl.h
@ -63,10 +63,6 @@ PairSNAPKokkos<DeviceType, real_type, vector_length>::PairSNAPKokkos(LAMMPS *lmp
  datamask_read = EMPTY_MASK;
  datamask_modify = EMPTY_MASK;

-  k_cutsq = tdual_fparams("PairSNAPKokkos::cutsq",atom->ntypes+1,atom->ntypes+1);
-  auto d_cutsq = k_cutsq.template view<DeviceType>();
-  rnd_cutsq = d_cutsq;
-
  host_flag = (execution_space == Host);
 }

@ -546,6 +542,9 @@ void PairSNAPKokkos<DeviceType, real_type, vector_length>::allocate()

  int n = atom->ntypes;
  MemKK::realloc_kokkos(d_map,"PairSNAPKokkos::map",n+1);
+
+  MemKK::realloc_kokkos(k_cutsq,"PairSNAPKokkos::cutsq",n+1,n+1);
+  rnd_cutsq = k_cutsq.template view<DeviceType>();
 }


--- a/src/KOKKOS/pair_table_kokkos.cpp
+++ b/src/KOKKOS/pair_table_kokkos.cpp
@ -133,19 +133,19 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
  EV_FLOAT ev;
  if (atom->ntypes > MAX_TYPES_STACKPARAMS) {
    if (neighflag == FULL) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,FULL,false,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,FULL,false,0,S_TableCompute<DeviceType,TABSTYLE> >
        ff(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,ff,ev);
      else Kokkos::parallel_for(list->inum,ff);
      ff.contribute();
    } else if (neighflag == HALFTHREAD) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,HALFTHREAD,false,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,HALFTHREAD,false,0,S_TableCompute<DeviceType,TABSTYLE> >
        ff(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,ff,ev);
      else Kokkos::parallel_for(list->inum,ff);
      ff.contribute();
    } else if (neighflag == HALF) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,HALF,false,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,HALF,false,0,S_TableCompute<DeviceType,TABSTYLE> >
        f(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,f,ev);
      else Kokkos::parallel_for(list->inum,f);
@ -153,19 +153,19 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in)
    }
  } else {
    if (neighflag == FULL) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,FULL,true,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,FULL,true,0,S_TableCompute<DeviceType,TABSTYLE> >
        f(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,f,ev);
      else Kokkos::parallel_for(list->inum,f);
      f.contribute();
    } else if (neighflag == HALFTHREAD) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,HALFTHREAD,true,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,HALFTHREAD,true,0,S_TableCompute<DeviceType,TABSTYLE> >
        f(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,f,ev);
      else Kokkos::parallel_for(list->inum,f);
      f.contribute();
    } else if (neighflag == HALF) {
-      PairComputeFunctor<PairTableKokkos<DeviceType>,HALF,true,S_TableCompute<DeviceType,TABSTYLE> >
+      PairComputeFunctor<PairTableKokkos<DeviceType>,HALF,true,0,S_TableCompute<DeviceType,TABSTYLE> >
        f(this,(NeighListKokkos<DeviceType>*) list);
      if (eflag || vflag) Kokkos::parallel_reduce(list->inum,f,ev);
      else Kokkos::parallel_for(list->inum,f);
--- a/src/KOKKOS/pair_table_kokkos.h
+++ b/src/KOKKOS/pair_table_kokkos.h
@ -35,9 +35,6 @@ struct S_TableCompute {
  static constexpr int TabStyle = TABSTYLE;
 };

-template <class DeviceType, int NEIGHFLAG, int TABSTYLE>
-struct PairTableComputeFunctor;
-
 template<class DeviceType>
 class PairTableKokkos : public PairTable {
 public:
@ -135,33 +132,33 @@ class PairTableKokkos : public PairTable {
  F_FLOAT compute_ecoul(const F_FLOAT& /*rsq*/, const int& /*i*/, const int& /*j*/,
                        const int& /*itype*/, const int& /*jtype*/) const { return 0; }

-  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,S_TableCompute<DeviceType,LOOKUP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,S_TableCompute<DeviceType,LOOKUP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,S_TableCompute<DeviceType,LOOKUP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,S_TableCompute<DeviceType,LOOKUP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,S_TableCompute<DeviceType,LOOKUP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,0,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,0,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,0,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,0,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,0,S_TableCompute<DeviceType,LOOKUP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,0,S_TableCompute<DeviceType,LOOKUP> >;

-  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,S_TableCompute<DeviceType,LINEAR> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,S_TableCompute<DeviceType,LINEAR> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,S_TableCompute<DeviceType,LINEAR> >;
-  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,S_TableCompute<DeviceType,LINEAR> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,S_TableCompute<DeviceType,LINEAR> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,0,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,0,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,0,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,0,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,0,S_TableCompute<DeviceType,LINEAR> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,0,S_TableCompute<DeviceType,LINEAR> >;

-  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,S_TableCompute<DeviceType,SPLINE> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,S_TableCompute<DeviceType,SPLINE> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,S_TableCompute<DeviceType,SPLINE> >;
-  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,S_TableCompute<DeviceType,SPLINE> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,S_TableCompute<DeviceType,SPLINE> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,0,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,0,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,0,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,0,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,0,S_TableCompute<DeviceType,SPLINE> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,0,S_TableCompute<DeviceType,SPLINE> >;

-  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,S_TableCompute<DeviceType,BITMAP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,S_TableCompute<DeviceType,BITMAP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,S_TableCompute<DeviceType,BITMAP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,S_TableCompute<DeviceType,BITMAP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,S_TableCompute<DeviceType,BITMAP> >;
-  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,true,0,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,true,0,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,true,0,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,FULL,false,0,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALF,false,0,S_TableCompute<DeviceType,BITMAP> >;
+  friend struct PairComputeFunctor<PairTableKokkos,HALFTHREAD,false,0,S_TableCompute<DeviceType,BITMAP> >;

  friend void pair_virial_fdotr_compute<PairTableKokkos>(PairTableKokkos*);
 };
--- a/src/KOKKOS/pair_yukawa_kokkos.h
+++ b/src/KOKKOS/pair_yukawa_kokkos.h
@ -95,20 +95,19 @@ class PairYukawaKokkos : public PairYukawa {
  int nlocal,nall,eflag,vflag;

  void allocate() override;
-  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairYukawaKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairYukawaKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairYukawaKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairYukawaKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairYukawaKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,FULL,void>(
-    PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,HALF,void>(
-    PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,HALFTHREAD,void>(
-    PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairYukawaKokkos,void>(
-    PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,FULL,0>(PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,FULL,1>(PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,HALF>(PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairYukawaKokkos,HALFTHREAD>(PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairYukawaKokkos,void>(PairYukawaKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairYukawaKokkos>(PairYukawaKokkos*);

 };
--- a/src/KOKKOS/pair_zbl_kokkos.h
+++ b/src/KOKKOS/pair_zbl_kokkos.h
@ -89,16 +89,19 @@ class PairZBLKokkos : public PairZBL {

  void allocate() override;

-  friend struct PairComputeFunctor<PairZBLKokkos,FULL,true>;
+  friend struct PairComputeFunctor<PairZBLKokkos,FULL,true,0>;
+  friend struct PairComputeFunctor<PairZBLKokkos,FULL,true,1>;
  friend struct PairComputeFunctor<PairZBLKokkos,HALF,true>;
  friend struct PairComputeFunctor<PairZBLKokkos,HALFTHREAD,true>;
-  friend struct PairComputeFunctor<PairZBLKokkos,FULL,false>;
+  friend struct PairComputeFunctor<PairZBLKokkos,FULL,false,0>;
+  friend struct PairComputeFunctor<PairZBLKokkos,FULL,false,1>;
  friend struct PairComputeFunctor<PairZBLKokkos,HALF,false>;
  friend struct PairComputeFunctor<PairZBLKokkos,HALFTHREAD,false>;
-  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,FULL,void>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,HALF,void>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,HALFTHREAD,void>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
-  friend EV_FLOAT pair_compute<PairZBLKokkos,void>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,FULL,0>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,FULL,1>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,HALF>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute_neighlist<PairZBLKokkos,HALFTHREAD>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
+  friend EV_FLOAT pair_compute<PairZBLKokkos>(PairZBLKokkos*,NeighListKokkos<DeviceType>*);
  friend void pair_virial_fdotr_compute<PairZBLKokkos>(PairZBLKokkos*);
 };

--- a/src/MAKE/OPTIONS/Makefile.intel_coprocessor
+++ b/src/MAKE/OPTIONS/Makefile.intel_coprocessor
@ -14,6 +14,7 @@ CCFLAGS =	-qopenmp -qoffload -ansi-alias -restrict \
                -DLMP_INTEL_USELRT -DLMP_USE_MKL_RNG -DLMP_INTEL_OFFLOAD \
                $(OPTFLAGS) -I$(MKLROOT)/include 
 SHFLAGS =	-fPIC
+FMTFLAGS = 	-std=c++11
 DEPFLAGS =	-M

 LINK =		mpiicpc -std=c++11
@ -118,6 +119,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.intel_cpu_intelmpi
+++ b/src/MAKE/OPTIONS/Makefile.intel_cpu_intelmpi
@ -6,13 +6,14 @@ SHELL = /bin/sh
 # compiler/linker settings
 # specify flags and libraries needed for your compiler

-CC =		mpiicpc -std=c++11 -diag-disable=10441 -diag-disable=2196
+CC =		mpiicpc -std=c++11 -diag-disable=10441 -diag-disable=2196 -diag-disable=11074 -diag-disable=11076
 OPTFLAGS =      -xHost -O2 -fp-model fast=2 -no-prec-div -qoverride-limits \
                -qopt-zmm-usage=high
 CCFLAGS =	-qopenmp -qno-offload -ansi-alias -restrict \
                -DLMP_INTEL_USELRT -DLMP_USE_MKL_RNG $(OPTFLAGS) \
                -I$(MKLROOT)/include
 SHFLAGS =	-fPIC
+FMTFLAGS = 	-std=c++11
 DEPFLAGS =	-M

 LINK =		mpiicpc -std=c++11 -diag-disable=10441 -diag-disable=2196
@ -117,6 +118,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.intel_cpu_mpich
+++ b/src/MAKE/OPTIONS/Makefile.intel_cpu_mpich
@ -6,13 +6,14 @@ SHELL = /bin/sh
 # compiler/linker settings
 # specify flags and libraries needed for your compiler

-CC =		mpicxx -cxx=icc -std=c++11 -diag-disable=10441 -diag-disable=2196
+CC =		mpicxx -cxx=icc -std=c++11 -diag-disable=10441 -diag-disable=2196 -diag-disable=11074 -diag-disable=11076
 OPTFLAGS =      -xHost -O2 -fp-model fast=2 -no-prec-div -qoverride-limits \
                -qopt-zmm-usage=high
 CCFLAGS =	-qopenmp -qno-offload -ansi-alias -restrict \
                -DLMP_INTEL_USELRT -DLMP_USE_MKL_RNG $(OPTFLAGS) \
                -I$(MKLROOT)/include
 SHFLAGS =	-fPIC
+FMTFLAGS = 	-std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx -cxx=icc -std=c++11 -diag-disable=10441 -diag-disable=2196
@ -117,6 +118,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.intel_cpu_openmpi
+++ b/src/MAKE/OPTIONS/Makefile.intel_cpu_openmpi
@ -7,13 +7,14 @@ SHELL = /bin/sh
 # specify flags and libraries needed for your compiler

 export OMPI_CXX = icc
-CC =		mpicxx -std=c++11 -diag-disable=10441 -diag-disable=2196
+CC =		mpicxx -std=c++11 -diag-disable=10441 -diag-disable=2196 -diag-disable=11074 -diag-disable=11076
 OPTFLAGS =      -xHost -O2 -fp-model fast=2 -no-prec-div -qoverride-limits \
                -qopt-zmm-usage=high
 CCFLAGS =	-qopenmp -qno-offload -ansi-alias -restrict \
                -DLMP_INTEL_USELRT -DLMP_USE_MKL_RNG $(OPTFLAGS) \
                -I$(MKLROOT)/include
 SHFLAGS =	-fPIC
+FMTFLAGS = 	-std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx -std=c++11 -diag-disable=10441 -diag-disable=2196
@ -118,6 +119,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.kokkos_cuda_mpi
+++ b/src/MAKE/OPTIONS/Makefile.kokkos_cuda_mpi
@ -12,6 +12,8 @@ export OMPI_CXX = $(KOKKOS_ABSOLUTE_PATH)/bin/nvcc_wrapper
 CC =		mpicxx
 CCFLAGS =	-g -O3 -DNDEBUG -Xcudafe --diag_suppress=unrecognized_pragma
 SHFLAGS =	-fPIC
+# uncomment when compiling with Intel 21.5 or older
+FMTFLAGS = 	# -std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx
@ -36,7 +38,7 @@ KOKKOS_ARCH = Volta70
 LMP_INC =	-DLAMMPS_GZIP

 # MPI library
-# see discussion in Section 2.2 (step 5) of manual
+# see discussion in Section 3.4 of the manual
 # MPI wrapper compiler/linker can provide this info
 # can point to dummy MPI library in src/STUBS as in Makefile.serial
 # use -D MPICH and OMPI settings in INC to avoid C++ lib conflicts
@ -118,6 +120,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.kokkos_mpi_only
+++ b/src/MAKE/OPTIONS/Makefile.kokkos_mpi_only
@ -9,6 +9,8 @@ SHELL = /bin/sh
 CC =		mpicxx
 CCFLAGS =	-g -O3 -DNDEBUG
 SHFLAGS =	-fPIC
+# uncomment when compiling with Intel 21.5 or older
+FMTFLAGS = 	# -std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx
@ -114,6 +116,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.kokkos_omp
+++ b/src/MAKE/OPTIONS/Makefile.kokkos_omp
@ -9,6 +9,8 @@ SHELL = /bin/sh
 CC =		mpicxx
 CCFLAGS =	-g -O3 -DNDEBUG
 SHFLAGS =	-fPIC
+# uncomment when compiling with Intel 21.5 or older
+FMTFLAGS = 	# -std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx
@ -114,6 +116,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/src/MAKE/OPTIONS/Makefile.kokkos_phi
+++ b/src/MAKE/OPTIONS/Makefile.kokkos_phi
@ -9,6 +9,8 @@ SHELL = /bin/sh
 CC =		mpicxx
 CCFLAGS =	-g -O3 -DNDEBUG
 SHFLAGS =	-fPIC
+# uncomment when compiling with Intel 21.5 or older
+FMTFLAGS = 	# -std=c++11
 DEPFLAGS =	-M

 LINK =		mpicxx
@ -115,6 +117,18 @@ $(SHLIB): $(OBJ) $(EXTRA_LINK_DEPENDS)
 %.o:%.cpp
 	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) -c $<

+variable.o : ../variable.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+utils.o : ../utils.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_format.o : ../fmtlib_format.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
+fmtlib_os.o : ../fmtlib_os.cpp
+	$(CC) $(CCFLAGS) $(SHFLAGS) $(EXTRA_INC) $(FMTFLAGS) -c $<
+
 # Individual dependencies

 depend : fastdep.exe $(SRC)
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Axel Kohlmeyer	27e8d0f19c	Merge pull request #3933 from lammps/maintenance Second Set of Collected Bug Fixes and Maintenance Updates for 2 August 2023 Stable release	2023-12-14 21:09:30 -05:00
Axel Kohlmeyer	9befd421ca	workaround hack for macOS	2023-12-14 18:07:50 -05:00
Axel Kohlmeyer	b3e54549db	safely copy balance shift dimension string with proper termination	2023-12-14 17:32:20 -05:00
Axel Kohlmeyer	85393862af	fix typos	2023-12-14 16:48:25 -05:00
Axel Kohlmeyer	ac1db251cb	recover compilation	2023-12-14 16:24:22 -05:00
Axel Kohlmeyer	3f48d48eea	add missing dependency	2023-12-14 16:00:10 -05:00
Stan Gerald Moore	d9804d7590	Fix issues with sorting neigh list by cutoff distance	2023-12-14 15:46:32 -05:00
Stan Gerald Moore	4128d52e1c	Bugfix: port missed changes from #3846	2023-12-14 15:45:51 -05:00
Axel Kohlmeyer	2d961e76b3	flag update #2 to stable release	2023-12-13 00:32:59 -05:00
Richard Berger	016c9ef4b2	Use PyConfig to initialize Python	2023-12-13 00:30:49 -05:00
Axel Kohlmeyer	e69c65431f	silence preprocessor warning from leaking internal define in cython generated code	2023-12-13 00:29:42 -05:00
Axel Kohlmeyer	a40e9222aa	add valgrind suppressions for MPICH on Fedora 39	2023-12-13 00:29:05 -05:00
Axel Kohlmeyer	283e2103e3	update fix adapt/fep from fix adapt. only supports 2-d parameters for pair styles	2023-12-06 14:35:05 -05:00
Axel Kohlmeyer	2808e6fc52	fix typo	2023-12-06 06:58:55 -05:00
Axel Kohlmeyer	c742b20c5a	update Purge.list and avoid redundant checks	2023-12-03 23:27:53 -05:00
Axel Kohlmeyer	530f487dd7	must do region check only when region is active	2023-12-03 11:22:35 -05:00
Axel Kohlmeyer	ba8ca9258b	correct dpi to get proper image scaling in PDF output	2023-12-02 16:35:01 -05:00
Axel Kohlmeyer	cd21f67cc6	avoid copying over terminating null	2023-12-02 16:22:43 -05:00
Axel Kohlmeyer	07257595ff	use r_c consistently	2023-12-01 05:52:49 -05:00
Axel Kohlmeyer	413d485617	recreate compute xrd mesh image with reasonable dpi setting and used PNG format	2023-12-01 01:32:28 -05:00
Axel Kohlmeyer	8759a18437	handle thermo_modify energy yes correctly	2023-11-30 10:33:54 -05:00
Axel Kohlmeyer	f79e9a113f	error out when no per-type masses are set. warn if both per-type and per-atom masses are used.	2023-11-27 07:47:55 -05:00
Axel Kohlmeyer	c1fa89186a	correct fix mvv/* compatibility checks in DPD-MESO package	2023-11-26 10:31:50 -05:00
Axel Kohlmeyer	609f5ec64b	restore using nvcc_wrapper with kokkos-cude.cmake preset	2023-11-25 05:58:05 -05:00
Axel Kohlmeyer	38b79eeb9b	some compilers require a code block to follow OpenMP pragmas, even if empty.	2023-11-24 14:53:53 -05:00
Axel Kohlmeyer	7035249abd	remove redundant code and fix memory leaks	2023-11-24 05:06:53 -05:00
Axel Kohlmeyer	816d74d80c	make compatible with Kokkos 3.7	2023-11-23 14:25:05 -05:00
Axel Kohlmeyer	4926164050	report Kokkos library version and OpenMP standard version	2023-11-23 12:38:59 -05:00
Axel Kohlmeyer	a102d64a95	detect newer OpenMP standard versions	2023-11-23 12:38:46 -05:00
Axel Kohlmeyer	77db8e422a	add check and document that "scale yes" is not supported for scaling atomic parameters with fix adapt/fep	2023-11-23 12:38:27 -05:00
Axel Kohlmeyer	ee0c5dc121	Update CODEOWNERS for cmake	2023-11-21 15:48:40 -05:00
Axel Kohlmeyer	184f5a7f5e	copy intel C++17 compiler hack to Kokkos makefiles	2023-11-21 13:00:09 -05:00
Axel Kohlmeyer	162b9c3ff3	tweak intel compiler makefile for traditional build	2023-11-21 12:59:54 -05:00
Axel Kohlmeyer	4d06a9928f	reduce warnings when compiling with intel classic compilers	2023-11-21 12:58:57 -05:00
Axel Kohlmeyer	938682a751	lower the C++ standard to 14 for some files when compiling with intel classic compiler	2023-11-21 12:58:33 -05:00
Axel Kohlmeyer	00bccbf067	check if creating unix domain socket failed	2023-11-17 03:20:23 -05:00
Axel Kohlmeyer	67085517ff	avoid segfault on command errors in force style unit tests and print error mesage instead	2023-11-16 22:10:15 -05:00
Axel Kohlmeyer	3a2d94822a	throw error for illegal replication values	2023-11-15 08:03:38 -05:00
Stan Gerald Moore	c272e8f94f	Avoid integer division	2023-11-15 07:35:26 -05:00
Stan Gerald Moore	7f41eb6d9a	Need force_clear for atom_vec_spin_kokkos	2023-11-15 07:35:12 -05:00
Stan Gerald Moore	a716df7e59	Fix bug in Kokkos minimize + fix deform	2023-11-15 07:34:57 -05:00
Axel Kohlmeyer	08eae40f9a	backport enforce2d with fix rigid bugfix	2023-11-15 07:10:55 -05:00
yury-lysogorskiy	b6c031fd03	Update pair_pace_extrapolation.cpp BUGFIX: pair_pace_extrapolation: setup flag aceimpl->ace->compute_projections = true before computing extrapolation grade	2023-11-10 11:51:15 -05:00
Richard Berger	990c07a133	bugfix: correctly build argv when using Python interface	2023-11-10 11:47:04 -05:00
Richard Berger	4e94e697ec	bugfix: make copy of exename	2023-11-10 11:46:53 -05:00
Richard Berger	4526dccaca	Correctly build argv with nullptr at the end	2023-11-10 11:46:40 -05:00
Stan Gerald Moore	917606e40e	Forces are not modified	2023-11-02 17:46:19 -04:00
Stan Gerald Moore	acaae8a36f	Fix bug in fix_dt_reset_kokkos	2023-11-02 17:46:10 -04:00
Axel Kohlmeyer	28803ee78d	add code to avoid deadlock	2023-11-02 02:17:23 -04:00
Steve Plimpton	dd498fcbf8	add comm of ghost atom coords to compute cluster/atom and aggregate/atom	2023-11-02 02:16:42 -04:00
Axel Kohlmeyer	0f8af20d0b	limit the maximum number of iterations so the LAMMPS simulation will not stall	2023-10-27 20:33:44 -04:00
Axel Kohlmeyer	00ef4ca3f6	fix bug in not listing all not compiled-in styles	2023-10-27 11:10:31 -04:00
Axel Kohlmeyer	50fbe61616	Backport of PR #3954 to stable release	2023-10-26 20:43:59 -04:00
Axel Kohlmeyer	854c6d93e2	more checks for misformatted ReST roles	2023-10-26 05:07:45 -04:00
Stan Gerald Moore	e8e2c5f986	Fix harmless compiler warnings	2023-10-24 17:23:42 -04:00
Axel Kohlmeyer	cff21ce808	improve help and error messages	2023-10-24 10:41:10 -04:00
Axel Kohlmeyer	97c4875a08	add sanity check on path to LAMMPS python package folder	2023-10-24 10:41:01 -04:00
Axel Kohlmeyer	c9aedf9df8	make sure liblinalg is built before linking phana	2023-10-23 14:58:04 -04:00
Axel Kohlmeyer	723dc17d80	must initialize deleted pointers to null since the following commands may fail	2023-10-23 07:35:10 -04:00
Axel Kohlmeyer	c90f874a0d	avoid invalid escape warnings for regexp expressions with python 3.12	2023-10-22 20:01:55 -04:00
Yifan Li	4ed5243d9b	add the missing dividing by np in compute t_prim	2023-10-21 14:58:46 -04:00
Axel Kohlmeyer	71c7d143b7	fix logic bug	2023-10-20 07:01:48 -04:00
Stan Moore	e944140ff2	whitespace	2023-10-19 15:29:45 -04:00
Stan Moore	b54545d1a4	Fix bug in Kokkos SNAP on GPUs	2023-10-19 15:29:33 -04:00
Stan Gerald Moore	fc7119982b	whitespace	2023-10-19 12:53:13 -04:00
Maria-Lesniewski	9e45df19c1	Barostat fix - see lammps PR 879 and 942	2023-10-19 12:25:24 -04:00
Stan Moore	8bfec75568	Add more error checks to Kokkos minimize	2023-10-19 10:11:42 -04:00
Axel Kohlmeyer	0f948e98f2	quote strings with special characters in keyword lists	2023-10-19 10:11:29 -04:00
Axel Kohlmeyer	b9ce258935	Revert "make sure itag is initialized" This reverts commit `058f87e019`.	2023-10-18 09:32:44 -04:00
Axel Kohlmeyer	058f87e019	make sure itag is initialized	2023-10-18 09:24:31 -04:00
Axel Kohlmeyer	6c2e469f5d	copy-and-paste bugfix from @stanmoore1	2023-10-17 19:40:09 -04:00
Stan Gerald Moore	810e3e5fa5	Fix issues with trim lists	2023-10-16 13:57:28 -04:00
Axel Kohlmeyer	a5374997d2	Revert "avoid issue with neighbor list trimming when used as a hybrid substyle" This reverts commit `23691d4336`.	2023-10-16 13:55:53 -04:00
Axel Kohlmeyer	e65ed32ecd	Revert "disable neighbor list trimming by default for REBO pair styles for now" This reverts commit `2ba7059c00`.	2023-10-16 13:55:52 -04:00
Axel Kohlmeyer	d326327bd7	Revert "disable neighbor list trimming for all other pair styles requesting neighbors of ghosts" This reverts commit `aa1c901f94`.	2023-10-16 13:55:48 -04:00
Axel Kohlmeyer	aa1c901f94	disable neighbor list trimming for all other pair styles requesting neighbors of ghosts	2023-10-16 00:02:10 -04:00
Axel Kohlmeyer	2ba7059c00	disable neighbor list trimming by default for REBO pair styles for now	2023-10-15 23:44:57 -04:00
Axel Kohlmeyer	23691d4336	avoid issue with neighbor list trimming when used as a hybrid substyle	2023-10-15 23:44:34 -04:00
Axel Kohlmeyer	78adc1727a	backport KOKKOS package fixes from PR #3930 by @stanmoore1	2023-10-13 16:32:36 -04:00
Axel Kohlmeyer	9def610c08	update PACE library	2023-10-13 16:31:09 -04:00
Axel Kohlmeyer	a939e93a08	must re-initialized threads also for neigbor lists	2023-10-11 17:42:33 -04:00
Axel Kohlmeyer	308207d5f9	fix cut-n-paste error	2023-10-05 13:16:47 -04:00
Evangelos Voyiatzis	75d0d9be1d	Fixes #3925 in region_ellipsoid.cpp	2023-10-04 10:56:13 -04:00
Axel Kohlmeyer	2f71bc7886	step LAMMPS GUI patch level number to indicate included bugfixes	2023-10-04 08:55:07 -04:00
Axel Kohlmeyer	ddbdaaafdc	make threads handling consistent. address issue that threads could not be increased	2023-10-04 08:46:01 -04:00
Axel Kohlmeyer	8946995199	enforce threads are reset properly for /omp styles	2023-10-04 08:39:58 -04:00
Axel Kohlmeyer	d567fdae97	fix delete / delete[] mismatch	2023-10-04 08:37:53 -04:00
Axel Kohlmeyer	ed9bfb433f	avoid segfaults when accessing lammps_last_thermo()	2023-10-04 08:35:42 -04:00
Axel Kohlmeyer	f8493ed805	Recognize Windows 11 23H2	2023-09-27 18:03:09 -04:00
Axel Kohlmeyer	f3beb206c9	support old ReaxFF force field files without ovcorr entry in bonds section	2023-09-27 00:06:15 -04:00
Axel Kohlmeyer	b5480e4e1b	must also update CWD when saving a file, not only when loading	2023-09-25 08:55:22 -04:00
Axel Kohlmeyer	f634b25e31	apply clang-format	2023-09-25 08:11:55 -04:00
Axel Kohlmeyer	b21db641d9	check for compatible LAMMPS version when creating LAMMPS instance This check must be done at runtime, since the LAMMPS shared library may have been loaded dynamically and thus required library functions may not be present or missing features with too only a LAMMPS version.	2023-09-25 08:08:00 -04:00
Axel Kohlmeyer	6ba94d1619	flag as maintenance branch again	2023-09-23 12:55:10 -04:00