diff --git a/src/KOKKOS/kokkos_type.h b/src/KOKKOS/kokkos_type.h index cc096058e..5b53b8ed0 100644 --- a/src/KOKKOS/kokkos_type.h +++ b/src/KOKKOS/kokkos_type.h @@ -1,929 +1,937 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifndef LMP_LMPTYPE_KOKKOS_H #define LMP_LMPTYPE_KOKKOS_H #include "lmptype.h" #include #include #include #include #if defined(KOKKOS_HAVE_CXX11) #undef ISFINITE #define ISFINITE(x) std::isfinite(x) #endif // User-settable FFT precision // FFT_PRECISION = 1 is single-precision complex (4-byte real, 4-byte imag) // FFT_PRECISION = 2 is double-precision complex (8-byte real, 8-byte imag) #ifdef FFT_SINGLE #define FFT_PRECISION 1 #define MPI_FFT_SCALAR MPI_FLOAT typedef float FFT_SCALAR; #else #define FFT_PRECISION 2 #define MPI_FFT_SCALAR MPI_DOUBLE typedef double FFT_SCALAR; #endif #define MAX_TYPES_STACKPARAMS 12 #define NeighClusterSize 8 struct lmp_float3 { float x,y,z; KOKKOS_INLINE_FUNCTION lmp_float3():x(0.0f),y(0.0f),z(0.0f) {} KOKKOS_INLINE_FUNCTION void operator += (const lmp_float3& tmp) { x+=tmp.x; y+=tmp.y; z+=tmp.z; } KOKKOS_INLINE_FUNCTION void operator += (const lmp_float3& tmp) volatile { x+=tmp.x; y+=tmp.y; z+=tmp.z; } KOKKOS_INLINE_FUNCTION void operator = (const lmp_float3& tmp) { x=tmp.x; y=tmp.y; z=tmp.z; } KOKKOS_INLINE_FUNCTION void operator = (const lmp_float3& tmp) volatile { x=tmp.x; y=tmp.y; z=tmp.z; } }; struct lmp_double3 { double x,y,z; KOKKOS_INLINE_FUNCTION lmp_double3():x(0.0),y(0.0),z(0.0) {} KOKKOS_INLINE_FUNCTION void operator += (const lmp_double3& tmp) { x+=tmp.x; y+=tmp.y; z+=tmp.z; } KOKKOS_INLINE_FUNCTION void operator += (const lmp_double3& tmp) volatile { x+=tmp.x; y+=tmp.y; z+=tmp.z; } KOKKOS_INLINE_FUNCTION void operator = (const lmp_double3& tmp) { x=tmp.x; y=tmp.y; z=tmp.z; } KOKKOS_INLINE_FUNCTION void operator = (const lmp_double3& tmp) volatile { x=tmp.x; y=tmp.y; z=tmp.z; } }; #if !defined(__CUDACC__) && !defined(__VECTOR_TYPES_H__) struct double2 { double x, y; }; struct float2 { float x, y; }; struct float4 { float x, y, z, w; }; struct double4 { double x, y, z, w; }; #endif // set LMPHostype and LMPDeviceType from Kokkos Default Types typedef Kokkos::DefaultExecutionSpace LMPDeviceType; typedef Kokkos::HostSpace::execution_space LMPHostType; // set ExecutionSpace stuct with variable "space" template struct ExecutionSpaceFromDevice; template<> struct ExecutionSpaceFromDevice { static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Host; }; #ifdef KOKKOS_HAVE_CUDA template<> struct ExecutionSpaceFromDevice { static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Device; }; #endif // define precision // handle global precision, force, energy, positions, kspace separately #ifndef PRECISION #define PRECISION 2 #endif #if PRECISION==1 typedef float LMP_FLOAT; typedef float2 LMP_FLOAT2; typedef lmp_float3 LMP_FLOAT3; typedef float4 LMP_FLOAT4; #else typedef double LMP_FLOAT; typedef double2 LMP_FLOAT2; typedef lmp_double3 LMP_FLOAT3; typedef double4 LMP_FLOAT4; #endif #ifndef PREC_FORCE #define PREC_FORCE PRECISION #endif #if PREC_FORCE==1 typedef float F_FLOAT; typedef float2 F_FLOAT2; typedef lmp_float3 F_FLOAT3; typedef float4 F_FLOAT4; #else typedef double F_FLOAT; typedef double2 F_FLOAT2; typedef lmp_double3 F_FLOAT3; typedef double4 F_FLOAT4; #endif #ifndef PREC_ENERGY #define PREC_ENERGY PRECISION #endif #if PREC_ENERGY==1 typedef float E_FLOAT; typedef float2 E_FLOAT2; typedef float4 E_FLOAT4; #else typedef double E_FLOAT; typedef double2 E_FLOAT2; typedef double4 E_FLOAT4; #endif struct s_EV_FLOAT { E_FLOAT evdwl; E_FLOAT ecoul; E_FLOAT v[6]; KOKKOS_INLINE_FUNCTION s_EV_FLOAT() { evdwl = 0; ecoul = 0; v[0] = 0; v[1] = 0; v[2] = 0; v[3] = 0; v[4] = 0; v[5] = 0; } KOKKOS_INLINE_FUNCTION void operator+=(const s_EV_FLOAT &rhs) { evdwl += rhs.evdwl; ecoul += rhs.ecoul; v[0] += rhs.v[0]; v[1] += rhs.v[1]; v[2] += rhs.v[2]; v[3] += rhs.v[3]; v[4] += rhs.v[4]; v[5] += rhs.v[5]; } KOKKOS_INLINE_FUNCTION void operator+=(const volatile s_EV_FLOAT &rhs) volatile { evdwl += rhs.evdwl; ecoul += rhs.ecoul; v[0] += rhs.v[0]; v[1] += rhs.v[1]; v[2] += rhs.v[2]; v[3] += rhs.v[3]; v[4] += rhs.v[4]; v[5] += rhs.v[5]; } }; typedef struct s_EV_FLOAT EV_FLOAT; struct s_EV_FLOAT_REAX { E_FLOAT evdwl; E_FLOAT ecoul; E_FLOAT v[6]; E_FLOAT ereax[10]; KOKKOS_INLINE_FUNCTION s_EV_FLOAT_REAX() { evdwl = 0; ecoul = 0; v[0] = 0; v[1] = 0; v[2] = 0; v[3] = 0; v[4] = 0; v[5] = 0; ereax[0] = 0; ereax[1] = 0; ereax[2] = 0; ereax[3] = 0; ereax[4] = 0; ereax[5] = 0; ereax[6] = 0; ereax[7] = 0; ereax[8] = 0; } KOKKOS_INLINE_FUNCTION void operator+=(const s_EV_FLOAT_REAX &rhs) { evdwl += rhs.evdwl; ecoul += rhs.ecoul; v[0] += rhs.v[0]; v[1] += rhs.v[1]; v[2] += rhs.v[2]; v[3] += rhs.v[3]; v[4] += rhs.v[4]; v[5] += rhs.v[5]; ereax[0] += rhs.ereax[0]; ereax[1] += rhs.ereax[1]; ereax[2] += rhs.ereax[2]; ereax[3] += rhs.ereax[3]; ereax[4] += rhs.ereax[4]; ereax[5] += rhs.ereax[5]; ereax[6] += rhs.ereax[6]; ereax[7] += rhs.ereax[7]; ereax[8] += rhs.ereax[8]; } KOKKOS_INLINE_FUNCTION void operator+=(const volatile s_EV_FLOAT_REAX &rhs) volatile { evdwl += rhs.evdwl; ecoul += rhs.ecoul; v[0] += rhs.v[0]; v[1] += rhs.v[1]; v[2] += rhs.v[2]; v[3] += rhs.v[3]; v[4] += rhs.v[4]; v[5] += rhs.v[5]; ereax[0] += rhs.ereax[0]; ereax[1] += rhs.ereax[1]; ereax[2] += rhs.ereax[2]; ereax[3] += rhs.ereax[3]; ereax[4] += rhs.ereax[4]; ereax[5] += rhs.ereax[5]; ereax[6] += rhs.ereax[6]; ereax[7] += rhs.ereax[7]; ereax[8] += rhs.ereax[8]; } }; typedef struct s_EV_FLOAT_REAX EV_FLOAT_REAX; #ifndef PREC_POS #define PREC_POS PRECISION #endif #if PREC_POS==1 typedef float X_FLOAT; typedef float2 X_FLOAT2; typedef float4 X_FLOAT4; #else typedef double X_FLOAT; typedef double2 X_FLOAT2; typedef double4 X_FLOAT4; #endif #ifndef PREC_VELOCITIES #define PREC_VELOCITIES PRECISION #endif #if PREC_VELOCITIES==1 typedef float V_FLOAT; typedef float2 V_FLOAT2; typedef float4 V_FLOAT4; #else typedef double V_FLOAT; typedef double2 V_FLOAT2; typedef double4 V_FLOAT4; #endif #if PREC_KSPACE==1 typedef float K_FLOAT; typedef float2 K_FLOAT2; typedef float4 K_FLOAT4; #else typedef double K_FLOAT; typedef double2 K_FLOAT2; typedef double4 K_FLOAT4; #endif // ------------------------------------------------------------------------ // LAMMPS types template struct ArrayTypes; template <> struct ArrayTypes { // scalar types typedef Kokkos:: DualView tdual_int_scalar; typedef tdual_int_scalar::t_dev t_int_scalar; typedef tdual_int_scalar::t_dev_const t_int_scalar_const; typedef tdual_int_scalar::t_dev_um t_int_scalar_um; typedef tdual_int_scalar::t_dev_const_um t_int_scalar_const_um; typedef Kokkos:: DualView tdual_float_scalar; typedef tdual_float_scalar::t_dev t_float_scalar; typedef tdual_float_scalar::t_dev_const t_float_scalar_const; typedef tdual_float_scalar::t_dev_um t_float_scalar_um; typedef tdual_float_scalar::t_dev_const_um t_float_scalar_const_um; // generic array types typedef Kokkos:: DualView tdual_int_1d; typedef tdual_int_1d::t_dev t_int_1d; typedef tdual_int_1d::t_dev_const t_int_1d_const; typedef tdual_int_1d::t_dev_um t_int_1d_um; typedef tdual_int_1d::t_dev_const_um t_int_1d_const_um; typedef tdual_int_1d::t_dev_const_randomread t_int_1d_randomread; typedef Kokkos:: DualView tdual_int_1d_3; typedef tdual_int_1d_3::t_dev t_int_1d_3; typedef tdual_int_1d_3::t_dev_const t_int_1d_3_const; typedef tdual_int_1d_3::t_dev_um t_int_1d_3_um; typedef tdual_int_1d_3::t_dev_const_um t_int_1d_3_const_um; typedef tdual_int_1d_3::t_dev_const_randomread t_int_1d_3_randomread; typedef Kokkos:: DualView tdual_int_2d; typedef tdual_int_2d::t_dev t_int_2d; typedef tdual_int_2d::t_dev_const t_int_2d_const; typedef tdual_int_2d::t_dev_um t_int_2d_um; typedef tdual_int_2d::t_dev_const_um t_int_2d_const_um; typedef tdual_int_2d::t_dev_const_randomread t_int_2d_randomread; typedef Kokkos:: DualView tdual_int_2d_dl; typedef tdual_int_2d_dl::t_dev t_int_2d_dl; typedef tdual_int_2d_dl::t_dev_const t_int_2d_const_dl; typedef tdual_int_2d_dl::t_dev_um t_int_2d_um_dl; typedef tdual_int_2d_dl::t_dev_const_um t_int_2d_const_um_dl; typedef tdual_int_2d_dl::t_dev_const_randomread t_int_2d_randomread_dl; typedef Kokkos:: DualView tdual_tagint_1d; typedef tdual_tagint_1d::t_dev t_tagint_1d; typedef tdual_tagint_1d::t_dev_const t_tagint_1d_const; typedef tdual_tagint_1d::t_dev_um t_tagint_1d_um; typedef tdual_tagint_1d::t_dev_const_um t_tagint_1d_const_um; typedef tdual_tagint_1d::t_dev_const_randomread t_tagint_1d_randomread; typedef Kokkos:: DualView tdual_tagint_2d; typedef tdual_tagint_2d::t_dev t_tagint_2d; typedef tdual_tagint_2d::t_dev_const t_tagint_2d_const; typedef tdual_tagint_2d::t_dev_um t_tagint_2d_um; typedef tdual_tagint_2d::t_dev_const_um t_tagint_2d_const_um; typedef tdual_tagint_2d::t_dev_const_randomread t_tagint_2d_randomread; typedef Kokkos:: DualView tdual_imageint_1d; typedef tdual_imageint_1d::t_dev t_imageint_1d; typedef tdual_imageint_1d::t_dev_const t_imageint_1d_const; typedef tdual_imageint_1d::t_dev_um t_imageint_1d_um; typedef tdual_imageint_1d::t_dev_const_um t_imageint_1d_const_um; typedef tdual_imageint_1d::t_dev_const_randomread t_imageint_1d_randomread; typedef Kokkos:: DualView tdual_double_1d; typedef tdual_double_1d::t_dev t_double_1d; typedef tdual_double_1d::t_dev_const t_double_1d_const; typedef tdual_double_1d::t_dev_um t_double_1d_um; typedef tdual_double_1d::t_dev_const_um t_double_1d_const_um; typedef tdual_double_1d::t_dev_const_randomread t_double_1d_randomread; typedef Kokkos:: DualView tdual_double_2d; typedef tdual_double_2d::t_dev t_double_2d; typedef tdual_double_2d::t_dev_const t_double_2d_const; typedef tdual_double_2d::t_dev_um t_double_2d_um; typedef tdual_double_2d::t_dev_const_um t_double_2d_const_um; typedef tdual_double_2d::t_dev_const_randomread t_double_2d_randomread; // 1d float array n typedef Kokkos::DualView tdual_float_1d; typedef tdual_float_1d::t_dev t_float_1d; typedef tdual_float_1d::t_dev_const t_float_1d_const; typedef tdual_float_1d::t_dev_um t_float_1d_um; typedef tdual_float_1d::t_dev_const_um t_float_1d_const_um; typedef tdual_float_1d::t_dev_const_randomread t_float_1d_randomread; //2d float array n typedef Kokkos::DualView tdual_float_2d; typedef tdual_float_2d::t_dev t_float_2d; typedef tdual_float_2d::t_dev_const t_float_2d_const; typedef tdual_float_2d::t_dev_um t_float_2d_um; typedef tdual_float_2d::t_dev_const_um t_float_2d_const_um; typedef tdual_float_2d::t_dev_const_randomread t_float_2d_randomread; //Position Types //1d X_FLOAT array n typedef Kokkos::DualView tdual_xfloat_1d; typedef tdual_xfloat_1d::t_dev t_xfloat_1d; typedef tdual_xfloat_1d::t_dev_const t_xfloat_1d_const; typedef tdual_xfloat_1d::t_dev_um t_xfloat_1d_um; typedef tdual_xfloat_1d::t_dev_const_um t_xfloat_1d_const_um; typedef tdual_xfloat_1d::t_dev_const_randomread t_xfloat_1d_randomread; //2d X_FLOAT array n*m typedef Kokkos::DualView tdual_xfloat_2d; typedef tdual_xfloat_2d::t_dev t_xfloat_2d; typedef tdual_xfloat_2d::t_dev_const t_xfloat_2d_const; typedef tdual_xfloat_2d::t_dev_um t_xfloat_2d_um; typedef tdual_xfloat_2d::t_dev_const_um t_xfloat_2d_const_um; typedef tdual_xfloat_2d::t_dev_const_randomread t_xfloat_2d_randomread; //2d X_FLOAT array n*4 #ifdef LMP_KOKKOS_NO_LEGACY typedef Kokkos::DualView tdual_x_array; #else typedef Kokkos::DualView tdual_x_array; #endif typedef tdual_x_array::t_dev t_x_array; typedef tdual_x_array::t_dev_const t_x_array_const; typedef tdual_x_array::t_dev_um t_x_array_um; typedef tdual_x_array::t_dev_const_um t_x_array_const_um; typedef tdual_x_array::t_dev_const_randomread t_x_array_randomread; //Velocity Types //1d V_FLOAT array n typedef Kokkos::DualView tdual_vfloat_1d; typedef tdual_vfloat_1d::t_dev t_vfloat_1d; typedef tdual_vfloat_1d::t_dev_const t_vfloat_1d_const; typedef tdual_vfloat_1d::t_dev_um t_vfloat_1d_um; typedef tdual_vfloat_1d::t_dev_const_um t_vfloat_1d_const_um; typedef tdual_vfloat_1d::t_dev_const_randomread t_vfloat_1d_randomread; //2d V_FLOAT array n*m typedef Kokkos::DualView tdual_vfloat_2d; typedef tdual_vfloat_2d::t_dev t_vfloat_2d; typedef tdual_vfloat_2d::t_dev_const t_vfloat_2d_const; typedef tdual_vfloat_2d::t_dev_um t_vfloat_2d_um; typedef tdual_vfloat_2d::t_dev_const_um t_vfloat_2d_const_um; typedef tdual_vfloat_2d::t_dev_const_randomread t_vfloat_2d_randomread; //2d V_FLOAT array n*3 typedef Kokkos::DualView tdual_v_array; //typedef Kokkos::DualView tdual_v_array; typedef tdual_v_array::t_dev t_v_array; typedef tdual_v_array::t_dev_const t_v_array_const; typedef tdual_v_array::t_dev_um t_v_array_um; typedef tdual_v_array::t_dev_const_um t_v_array_const_um; typedef tdual_v_array::t_dev_const_randomread t_v_array_randomread; //Force Types //1d F_FLOAT array n typedef Kokkos::DualView tdual_ffloat_1d; typedef tdual_ffloat_1d::t_dev t_ffloat_1d; typedef tdual_ffloat_1d::t_dev_const t_ffloat_1d_const; typedef tdual_ffloat_1d::t_dev_um t_ffloat_1d_um; typedef tdual_ffloat_1d::t_dev_const_um t_ffloat_1d_const_um; typedef tdual_ffloat_1d::t_dev_const_randomread t_ffloat_1d_randomread; //2d F_FLOAT array n*m typedef Kokkos::DualView tdual_ffloat_2d; typedef tdual_ffloat_2d::t_dev t_ffloat_2d; typedef tdual_ffloat_2d::t_dev_const t_ffloat_2d_const; typedef tdual_ffloat_2d::t_dev_um t_ffloat_2d_um; typedef tdual_ffloat_2d::t_dev_const_um t_ffloat_2d_const_um; typedef tdual_ffloat_2d::t_dev_const_randomread t_ffloat_2d_randomread; //2d F_FLOAT array n*m, device layout typedef Kokkos::DualView tdual_ffloat_2d_dl; typedef tdual_ffloat_2d_dl::t_dev t_ffloat_2d_dl; typedef tdual_ffloat_2d_dl::t_dev_const t_ffloat_2d_const_dl; typedef tdual_ffloat_2d_dl::t_dev_um t_ffloat_2d_um_dl; typedef tdual_ffloat_2d_dl::t_dev_const_um t_ffloat_2d_const_um_dl; typedef tdual_ffloat_2d_dl::t_dev_const_randomread t_ffloat_2d_randomread_dl; //2d F_FLOAT array n*3 typedef Kokkos::DualView tdual_f_array; //typedef Kokkos::DualView tdual_f_array; typedef tdual_f_array::t_dev t_f_array; typedef tdual_f_array::t_dev_const t_f_array_const; typedef tdual_f_array::t_dev_um t_f_array_um; typedef tdual_f_array::t_dev_const_um t_f_array_const_um; typedef tdual_f_array::t_dev_const_randomread t_f_array_randomread; //2d F_FLOAT array n*6 (for virial) typedef Kokkos::DualView tdual_virial_array; typedef tdual_virial_array::t_dev t_virial_array; typedef tdual_virial_array::t_dev_const t_virial_array_const; typedef tdual_virial_array::t_dev_um t_virial_array_um; typedef tdual_virial_array::t_dev_const_um t_virial_array_const_um; typedef tdual_virial_array::t_dev_const_randomread t_virial_array_randomread; //Energy Types //1d E_FLOAT array n typedef Kokkos::DualView tdual_efloat_1d; typedef tdual_efloat_1d::t_dev t_efloat_1d; typedef tdual_efloat_1d::t_dev_const t_efloat_1d_const; typedef tdual_efloat_1d::t_dev_um t_efloat_1d_um; typedef tdual_efloat_1d::t_dev_const_um t_efloat_1d_const_um; typedef tdual_efloat_1d::t_dev_const_randomread t_efloat_1d_randomread; //2d E_FLOAT array n*m typedef Kokkos::DualView tdual_efloat_2d; typedef tdual_efloat_2d::t_dev t_efloat_2d; typedef tdual_efloat_2d::t_dev_const t_efloat_2d_const; typedef tdual_efloat_2d::t_dev_um t_efloat_2d_um; typedef tdual_efloat_2d::t_dev_const_um t_efloat_2d_const_um; typedef tdual_efloat_2d::t_dev_const_randomread t_efloat_2d_randomread; //2d E_FLOAT array n*3 typedef Kokkos::DualView tdual_e_array; typedef tdual_e_array::t_dev t_e_array; typedef tdual_e_array::t_dev_const t_e_array_const; typedef tdual_e_array::t_dev_um t_e_array_um; typedef tdual_e_array::t_dev_const_um t_e_array_const_um; typedef tdual_e_array::t_dev_const_randomread t_e_array_randomread; //Neighbor Types typedef Kokkos::DualView tdual_neighbors_2d; typedef tdual_neighbors_2d::t_dev t_neighbors_2d; typedef tdual_neighbors_2d::t_dev_const t_neighbors_2d_const; typedef tdual_neighbors_2d::t_dev_um t_neighbors_2d_um; typedef tdual_neighbors_2d::t_dev_const_um t_neighbors_2d_const_um; typedef tdual_neighbors_2d::t_dev_const_randomread t_neighbors_2d_randomread; //Kspace typedef Kokkos:: DualView tdual_FFT_SCALAR_1d; typedef tdual_FFT_SCALAR_1d::t_dev t_FFT_SCALAR_1d; typedef tdual_FFT_SCALAR_1d::t_dev_um t_FFT_SCALAR_1d_um; typedef Kokkos::DualView tdual_FFT_SCALAR_2d; typedef tdual_FFT_SCALAR_2d::t_dev t_FFT_SCALAR_2d; typedef Kokkos::DualView tdual_FFT_SCALAR_2d_3; typedef tdual_FFT_SCALAR_2d_3::t_dev t_FFT_SCALAR_2d_3; typedef Kokkos::DualView tdual_FFT_SCALAR_3d; typedef tdual_FFT_SCALAR_3d::t_dev t_FFT_SCALAR_3d; typedef Kokkos:: DualView tdual_FFT_DATA_1d; typedef tdual_FFT_DATA_1d::t_dev t_FFT_DATA_1d; typedef tdual_FFT_DATA_1d::t_dev_um t_FFT_DATA_1d_um; typedef Kokkos:: DualView tdual_int_64; typedef tdual_int_64::t_dev t_int_64; typedef tdual_int_64::t_dev_um t_int_64_um; }; #ifdef KOKKOS_HAVE_CUDA template <> struct ArrayTypes { //Scalar Types typedef Kokkos::DualView tdual_int_scalar; typedef tdual_int_scalar::t_host t_int_scalar; typedef tdual_int_scalar::t_host_const t_int_scalar_const; typedef tdual_int_scalar::t_host_um t_int_scalar_um; typedef tdual_int_scalar::t_host_const_um t_int_scalar_const_um; typedef Kokkos::DualView tdual_float_scalar; typedef tdual_float_scalar::t_host t_float_scalar; typedef tdual_float_scalar::t_host_const t_float_scalar_const; typedef tdual_float_scalar::t_host_um t_float_scalar_um; typedef tdual_float_scalar::t_host_const_um t_float_scalar_const_um; //Generic ArrayTypes typedef Kokkos::DualView tdual_int_1d; typedef tdual_int_1d::t_host t_int_1d; typedef tdual_int_1d::t_host_const t_int_1d_const; typedef tdual_int_1d::t_host_um t_int_1d_um; typedef tdual_int_1d::t_host_const_um t_int_1d_const_um; typedef tdual_int_1d::t_host_const_randomread t_int_1d_randomread; typedef Kokkos::DualView tdual_int_1d_3; typedef tdual_int_1d_3::t_host t_int_1d_3; typedef tdual_int_1d_3::t_host_const t_int_1d_3_const; typedef tdual_int_1d_3::t_host_um t_int_1d_3_um; typedef tdual_int_1d_3::t_host_const_um t_int_1d_3_const_um; typedef tdual_int_1d_3::t_host_const_randomread t_int_1d_3_randomread; typedef Kokkos::DualView tdual_int_2d; typedef tdual_int_2d::t_host t_int_2d; typedef tdual_int_2d::t_host_const t_int_2d_const; typedef tdual_int_2d::t_host_um t_int_2d_um; typedef tdual_int_2d::t_host_const_um t_int_2d_const_um; typedef tdual_int_2d::t_host_const_randomread t_int_2d_randomread; typedef Kokkos::DualView tdual_int_2d_dl; typedef tdual_int_2d_dl::t_host t_int_2d_dl; typedef tdual_int_2d_dl::t_host_const t_int_2d_const_dl; typedef tdual_int_2d_dl::t_host_um t_int_2d_um_dl; typedef tdual_int_2d_dl::t_host_const_um t_int_2d_const_um_dl; typedef tdual_int_2d_dl::t_host_const_randomread t_int_2d_randomread_dl; typedef Kokkos::DualView tdual_tagint_1d; typedef tdual_tagint_1d::t_host t_tagint_1d; typedef tdual_tagint_1d::t_host_const t_tagint_1d_const; typedef tdual_tagint_1d::t_host_um t_tagint_1d_um; typedef tdual_tagint_1d::t_host_const_um t_tagint_1d_const_um; typedef tdual_tagint_1d::t_host_const_randomread t_tagint_1d_randomread; typedef Kokkos:: DualView tdual_tagint_2d; typedef tdual_tagint_2d::t_host t_tagint_2d; typedef tdual_tagint_2d::t_host_const t_tagint_2d_const; typedef tdual_tagint_2d::t_host_um t_tagint_2d_um; typedef tdual_tagint_2d::t_host_const_um t_tagint_2d_const_um; typedef tdual_tagint_2d::t_host_const_randomread t_tagint_2d_randomread; typedef Kokkos:: DualView tdual_imageint_1d; typedef tdual_imageint_1d::t_host t_imageint_1d; typedef tdual_imageint_1d::t_host_const t_imageint_1d_const; typedef tdual_imageint_1d::t_host_um t_imageint_1d_um; typedef tdual_imageint_1d::t_host_const_um t_imageint_1d_const_um; typedef tdual_imageint_1d::t_host_const_randomread t_imageint_1d_randomread; typedef Kokkos:: DualView tdual_double_1d; typedef tdual_double_1d::t_host t_double_1d; typedef tdual_double_1d::t_host_const t_double_1d_const; typedef tdual_double_1d::t_host_um t_double_1d_um; typedef tdual_double_1d::t_host_const_um t_double_1d_const_um; typedef tdual_double_1d::t_host_const_randomread t_double_1d_randomread; typedef Kokkos:: DualView tdual_double_2d; typedef tdual_double_2d::t_host t_double_2d; typedef tdual_double_2d::t_host_const t_double_2d_const; typedef tdual_double_2d::t_host_um t_double_2d_um; typedef tdual_double_2d::t_host_const_um t_double_2d_const_um; typedef tdual_double_2d::t_host_const_randomread t_double_2d_randomread; //1d float array n typedef Kokkos::DualView tdual_float_1d; typedef tdual_float_1d::t_host t_float_1d; typedef tdual_float_1d::t_host_const t_float_1d_const; typedef tdual_float_1d::t_host_um t_float_1d_um; typedef tdual_float_1d::t_host_const_um t_float_1d_const_um; typedef tdual_float_1d::t_host_const_randomread t_float_1d_randomread; //2d float array n typedef Kokkos::DualView tdual_float_2d; typedef tdual_float_2d::t_host t_float_2d; typedef tdual_float_2d::t_host_const t_float_2d_const; typedef tdual_float_2d::t_host_um t_float_2d_um; typedef tdual_float_2d::t_host_const_um t_float_2d_const_um; typedef tdual_float_2d::t_host_const_randomread t_float_2d_randomread; //Position Types //1d X_FLOAT array n typedef Kokkos::DualView tdual_xfloat_1d; typedef tdual_xfloat_1d::t_host t_xfloat_1d; typedef tdual_xfloat_1d::t_host_const t_xfloat_1d_const; typedef tdual_xfloat_1d::t_host_um t_xfloat_1d_um; typedef tdual_xfloat_1d::t_host_const_um t_xfloat_1d_const_um; typedef tdual_xfloat_1d::t_host_const_randomread t_xfloat_1d_randomread; //2d X_FLOAT array n*m typedef Kokkos::DualView tdual_xfloat_2d; typedef tdual_xfloat_2d::t_host t_xfloat_2d; typedef tdual_xfloat_2d::t_host_const t_xfloat_2d_const; typedef tdual_xfloat_2d::t_host_um t_xfloat_2d_um; typedef tdual_xfloat_2d::t_host_const_um t_xfloat_2d_const_um; typedef tdual_xfloat_2d::t_host_const_randomread t_xfloat_2d_randomread; //2d X_FLOAT array n*3 typedef Kokkos::DualView tdual_x_array; typedef tdual_x_array::t_host t_x_array; typedef tdual_x_array::t_host_const t_x_array_const; typedef tdual_x_array::t_host_um t_x_array_um; typedef tdual_x_array::t_host_const_um t_x_array_const_um; typedef tdual_x_array::t_host_const_randomread t_x_array_randomread; //Velocity Types //1d V_FLOAT array n typedef Kokkos::DualView tdual_vfloat_1d; typedef tdual_vfloat_1d::t_host t_vfloat_1d; typedef tdual_vfloat_1d::t_host_const t_vfloat_1d_const; typedef tdual_vfloat_1d::t_host_um t_vfloat_1d_um; typedef tdual_vfloat_1d::t_host_const_um t_vfloat_1d_const_um; typedef tdual_vfloat_1d::t_host_const_randomread t_vfloat_1d_randomread; //2d V_FLOAT array n*m typedef Kokkos::DualView tdual_vfloat_2d; typedef tdual_vfloat_2d::t_host t_vfloat_2d; typedef tdual_vfloat_2d::t_host_const t_vfloat_2d_const; typedef tdual_vfloat_2d::t_host_um t_vfloat_2d_um; typedef tdual_vfloat_2d::t_host_const_um t_vfloat_2d_const_um; typedef tdual_vfloat_2d::t_host_const_randomread t_vfloat_2d_randomread; //2d V_FLOAT array n*3 typedef Kokkos::DualView tdual_v_array; //typedef Kokkos::DualView tdual_v_array; typedef tdual_v_array::t_host t_v_array; typedef tdual_v_array::t_host_const t_v_array_const; typedef tdual_v_array::t_host_um t_v_array_um; typedef tdual_v_array::t_host_const_um t_v_array_const_um; typedef tdual_v_array::t_host_const_randomread t_v_array_randomread; //Force Types //1d F_FLOAT array n typedef Kokkos::DualView tdual_ffloat_1d; typedef tdual_ffloat_1d::t_host t_ffloat_1d; typedef tdual_ffloat_1d::t_host_const t_ffloat_1d_const; typedef tdual_ffloat_1d::t_host_um t_ffloat_1d_um; typedef tdual_ffloat_1d::t_host_const_um t_ffloat_1d_const_um; typedef tdual_ffloat_1d::t_host_const_randomread t_ffloat_1d_randomread; //2d F_FLOAT array n*m typedef Kokkos::DualView tdual_ffloat_2d; typedef tdual_ffloat_2d::t_host t_ffloat_2d; typedef tdual_ffloat_2d::t_host_const t_ffloat_2d_const; typedef tdual_ffloat_2d::t_host_um t_ffloat_2d_um; typedef tdual_ffloat_2d::t_host_const_um t_ffloat_2d_const_um; typedef tdual_ffloat_2d::t_host_const_randomread t_ffloat_2d_randomread; //2d F_FLOAT array n*m, device layout typedef Kokkos::DualView tdual_ffloat_2d_dl; typedef tdual_ffloat_2d_dl::t_host t_ffloat_2d_dl; typedef tdual_ffloat_2d_dl::t_host_const t_ffloat_2d_const_dl; typedef tdual_ffloat_2d_dl::t_host_um t_ffloat_2d_um_dl; typedef tdual_ffloat_2d_dl::t_host_const_um t_ffloat_2d_const_um_dl; typedef tdual_ffloat_2d_dl::t_host_const_randomread t_ffloat_2d_randomread_dl; //2d F_FLOAT array n*3 typedef Kokkos::DualView tdual_f_array; //typedef Kokkos::DualView tdual_f_array; typedef tdual_f_array::t_host t_f_array; typedef tdual_f_array::t_host_const t_f_array_const; typedef tdual_f_array::t_host_um t_f_array_um; typedef tdual_f_array::t_host_const_um t_f_array_const_um; typedef tdual_f_array::t_host_const_randomread t_f_array_randomread; //2d F_FLOAT array n*6 (for virial) typedef Kokkos::DualView tdual_virial_array; typedef tdual_virial_array::t_host t_virial_array; typedef tdual_virial_array::t_host_const t_virial_array_const; typedef tdual_virial_array::t_host_um t_virial_array_um; typedef tdual_virial_array::t_host_const_um t_virial_array_const_um; typedef tdual_virial_array::t_host_const_randomread t_virial_array_randomread; //Energy Types //1d E_FLOAT array n typedef Kokkos::DualView tdual_efloat_1d; typedef tdual_efloat_1d::t_host t_efloat_1d; typedef tdual_efloat_1d::t_host_const t_efloat_1d_const; typedef tdual_efloat_1d::t_host_um t_efloat_1d_um; typedef tdual_efloat_1d::t_host_const_um t_efloat_1d_const_um; typedef tdual_efloat_1d::t_host_const_randomread t_efloat_1d_randomread; //2d E_FLOAT array n*m typedef Kokkos::DualView tdual_efloat_2d; typedef tdual_efloat_2d::t_host t_efloat_2d; typedef tdual_efloat_2d::t_host_const t_efloat_2d_const; typedef tdual_efloat_2d::t_host_um t_efloat_2d_um; typedef tdual_efloat_2d::t_host_const_um t_efloat_2d_const_um; typedef tdual_efloat_2d::t_host_const_randomread t_efloat_2d_randomread; //2d E_FLOAT array n*3 typedef Kokkos::DualView tdual_e_array; typedef tdual_e_array::t_host t_e_array; typedef tdual_e_array::t_host_const t_e_array_const; typedef tdual_e_array::t_host_um t_e_array_um; typedef tdual_e_array::t_host_const_um t_e_array_const_um; typedef tdual_e_array::t_host_const_randomread t_e_array_randomread; //Neighbor Types typedef Kokkos::DualView tdual_neighbors_2d; typedef tdual_neighbors_2d::t_host t_neighbors_2d; typedef tdual_neighbors_2d::t_host_const t_neighbors_2d_const; typedef tdual_neighbors_2d::t_host_um t_neighbors_2d_um; typedef tdual_neighbors_2d::t_host_const_um t_neighbors_2d_const_um; typedef tdual_neighbors_2d::t_host_const_randomread t_neighbors_2d_randomread; //Kspace typedef Kokkos:: DualView tdual_FFT_SCALAR_1d; typedef tdual_FFT_SCALAR_1d::t_host t_FFT_SCALAR_1d; typedef tdual_FFT_SCALAR_1d::t_host_um t_FFT_SCALAR_1d_um; typedef Kokkos::DualView tdual_FFT_SCALAR_2d; typedef tdual_FFT_SCALAR_2d::t_host t_FFT_SCALAR_2d; typedef Kokkos::DualView tdual_FFT_SCALAR_2d_3; typedef tdual_FFT_SCALAR_2d_3::t_host t_FFT_SCALAR_2d_3; typedef Kokkos::DualView tdual_FFT_SCALAR_3d; typedef tdual_FFT_SCALAR_3d::t_host t_FFT_SCALAR_3d; typedef Kokkos:: DualView tdual_FFT_DATA_1d; typedef tdual_FFT_DATA_1d::t_host t_FFT_DATA_1d; typedef tdual_FFT_DATA_1d::t_host_um t_FFT_DATA_1d_um; typedef Kokkos:: DualView tdual_int_64; typedef tdual_int_64::t_host t_int_64; typedef tdual_int_64::t_host_um t_int_64_um; }; #endif //default LAMMPS Types typedef struct ArrayTypes DAT; typedef struct ArrayTypes HAT; template void buffer_view(BufferView &buf, DualView &view, const size_t n0, const size_t n1 = 0, const size_t n2 = 0, const size_t n3 = 0, const size_t n4 = 0, const size_t n5 = 0, const size_t n6 = 0, const size_t n7 = 0) { buf = BufferView( view.template view().ptr_on_device(), n0,n1,n2,n3,n4,n5,n6,n7); } template struct MemsetZeroFunctor { typedef DeviceType execution_space ; void* ptr; KOKKOS_INLINE_FUNCTION void operator()(const int i) const { ((int*)ptr)[i] = 0; } }; template void memset_kokkos (ViewType &view) { static MemsetZeroFunctor f; f.ptr = view.ptr_on_device(); #ifndef KOKKOS_USING_DEPRECATED_VIEW Kokkos::parallel_for(view.span()*sizeof(typename ViewType::value_type)/4, f); #else Kokkos::parallel_for(view.capacity()*sizeof(typename ViewType::value_type)/4, f); #endif ViewType::execution_space::fence(); } +struct params_lj_coul { + KOKKOS_INLINE_FUNCTION + params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; + KOKKOS_INLINE_FUNCTION + params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; + F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; +}; + #if defined(KOKKOS_HAVE_CXX11) #undef ISFINITE #define ISFINITE(x) std::isfinite(x) #endif #ifdef KOKKOS_HAVE_CUDA #define LAMMPS_LAMBDA [=] __device__ #else #define LAMMPS_LAMBDA [=] #endif #endif diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h index 3c0b7d46a..048a7dab6 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.h @@ -1,165 +1,160 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/charmm/coul/charmm/implicit/kk,PairLJCharmmCoulCharmmImplicitKokkos) PairStyle(lj/charmm/coul/charmm/implicit/kk/device,PairLJCharmmCoulCharmmImplicitKokkos) PairStyle(lj/charmm/coul/charmm/implicit/kk/host,PairLJCharmmCoulCharmmImplicitKokkos) #else #ifndef LMP_PAIR_LJ_CHARMM_COUL_CHARMM_IMPLICIT_KOKKOS_H #define LMP_PAIR_LJ_CHARMM_COUL_CHARMM_IMPLICIT_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_charmm_coul_charmm_implicit.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCharmmCoulCharmmImplicitKokkos : public PairLJCharmmCoulCharmmImplicit { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCharmmCoulCharmmImplicitKokkos(class LAMMPS *); ~PairLJCharmmCoulCharmmImplicitKokkos(); void compute(int, int); void settings(int, char **); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulCharmmImplicitKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmImplicitKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulCharmmImplicitKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJCharmmCoulCharmmImplicitKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/charmm/coul/charmm/implicit/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h index 202cda68b..db0b14a84 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.h @@ -1,165 +1,160 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/charmm/coul/charmm/kk,PairLJCharmmCoulCharmmKokkos) PairStyle(lj/charmm/coul/charmm/kk/device,PairLJCharmmCoulCharmmKokkos) PairStyle(lj/charmm/coul/charmm/kk/host,PairLJCharmmCoulCharmmKokkos) #else #ifndef LMP_PAIR_LJ_CHARMM_COUL_CHARMM_KOKKOS_H #define LMP_PAIR_LJ_CHARMM_COUL_CHARMM_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_charmm_coul_charmm.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCharmmCoulCharmmKokkos : public PairLJCharmmCoulCharmm { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCharmmCoulCharmmKokkos(class LAMMPS *); ~PairLJCharmmCoulCharmmKokkos(); void compute(int, int); void settings(int, char **); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulCharmmKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulCharmmKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulCharmmKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJCharmmCoulCharmmKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/charmm/coul/charmm/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h index fcdab7ddb..0969d11b0 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.h @@ -1,158 +1,152 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/charmm/coul/long/kk,PairLJCharmmCoulLongKokkos) PairStyle(lj/charmm/coul/long/kk/device,PairLJCharmmCoulLongKokkos) PairStyle(lj/charmm/coul/long/kk/host,PairLJCharmmCoulLongKokkos) #else #ifndef LMP_PAIR_LJ_CHARMM_COUL_LONG_KOKKOS_H #define LMP_PAIR_LJ_CHARMM_COUL_LONG_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_charmm_coul_long.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCharmmCoulLongKokkos : public PairLJCharmmCoulLong { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCharmmCoulLongKokkos(class LAMMPS *); ~PairLJCharmmCoulLongKokkos(); void compute(int, int); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulLongKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCharmmCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCharmmCoulLongKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJCharmmCoulLongKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/charmm/coul/long/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h index 1ea5bc69b..c3492666d 100644 --- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.h @@ -1,144 +1,139 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/class2/coul/cut/kk,PairLJClass2CoulCutKokkos) PairStyle(lj/class2/coul/cut/kk/device,PairLJClass2CoulCutKokkos) PairStyle(lj/class2/coul/cut/kk/host,PairLJClass2CoulCutKokkos) #else #ifndef LMP_PAIR_LJ_CLASS2_COUL_CUT_KOKKOS_H #define LMP_PAIR_LJ_CLASS2_COUL_CUT_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_class2_coul_cut.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJClass2CoulCutKokkos : public PairLJClass2CoulCut { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJClass2CoulCutKokkos(class LAMMPS *); ~PairLJClass2CoulCutKokkos(); void compute(int, int); void settings(int, char **); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend EV_FLOAT pair_compute_neighlist(PairLJClass2CoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJClass2CoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJClass2CoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJClass2CoulCutKokkos*, NeighListKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/class2/coul/cut/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h index 0b1b2dc90..c5c46ed2d 100644 --- a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.h @@ -1,161 +1,155 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/class2/coul/long/kk,PairLJClass2CoulLongKokkos) PairStyle(lj/class2/coul/long/kk/device,PairLJClass2CoulLongKokkos) PairStyle(lj/class2/coul/long/kk/host,PairLJClass2CoulLongKokkos) #else #ifndef LMP_PAIR_LJ_CLASS2_COUL_LONG_KOKKOS_H #define LMP_PAIR_LJ_CLASS2_COUL_LONG_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_class2_coul_long.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJClass2CoulLongKokkos : public PairLJClass2CoulLong { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJClass2CoulLongKokkos(class LAMMPS *); ~PairLJClass2CoulLongKokkos(); void compute(int, int); void settings(int, char **); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJClass2CoulLongKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJClass2CoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJClass2CoulLongKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJClass2CoulLongKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/class2/coul/long/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h index 36f31d176..5891371d1 100644 --- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.h @@ -1,147 +1,139 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/cut/coul/cut/kk,PairLJCutCoulCutKokkos) PairStyle(lj/cut/coul/cut/kk/device,PairLJCutCoulCutKokkos) PairStyle(lj/cut/coul/cut/kk/host,PairLJCutCoulCutKokkos) #else #ifndef LMP_PAIR_LJ_CUT_COUL_CUT_KOKKOS_H #define LMP_PAIR_LJ_CUT_COUL_CUT_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_cut_coul_cut.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCutCoulCutKokkos : public PairLJCutCoulCut { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCutCoulCutKokkos(class LAMMPS *); ~PairLJCutCoulCutKokkos(); void compute(int, int); void settings(int, char **); void init_style(); double init_one(int, int); - struct params_lj_coul{ - KOKKOS_INLINE_FUNCTION - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - KOKKOS_INLINE_FUNCTION - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulCutKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulCutKokkos*, NeighListKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/cut/coul/cut/kk That style is not supported by Kokkos. */ diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h index 9e1e30aba..d507f76a3 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.h @@ -1,145 +1,139 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/cut/coul/debye/kk,PairLJCutCoulDebyeKokkos) PairStyle(lj/cut/coul/debye/kk/device,PairLJCutCoulDebyeKokkos) PairStyle(lj/cut/coul/debye/kk/host,PairLJCutCoulDebyeKokkos) #else #ifndef LMP_PAIR_LJ_CUT_COUL_DEBYE_KOKKOS_H #define LMP_PAIR_LJ_CUT_COUL_DEBYE_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_cut_coul_debye.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCutCoulDebyeKokkos : public PairLJCutCoulDebye { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCutCoulDebyeKokkos(class LAMMPS *); ~PairLJCutCoulDebyeKokkos(); void compute(int, int); void settings(int, char **); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDebyeKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDebyeKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDebyeKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulDebyeKokkos*, NeighListKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/cut/coul/debye/kk Self-explanatory. */ diff --git a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h index b1f578ec0..3e378757c 100644 --- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.h @@ -1,138 +1,132 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/cut/coul/dsf/kk,PairLJCutCoulDSFKokkos) PairStyle(lj/cut/coul/dsf/kk/device,PairLJCutCoulDSFKokkos) PairStyle(lj/cut/coul/dsf/kk/host,PairLJCutCoulDSFKokkos) #else #ifndef LMP_PAIR_LJ_CUT_COUL_DSF_KOKKOS_H #define LMP_PAIR_LJ_CUT_COUL_DSF_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_cut_coul_dsf.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCutCoulDSFKokkos : public PairLJCutCoulDSF { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCutCoulDSFKokkos(class LAMMPS *); ~PairLJCutCoulDSFKokkos(); void compute(int, int); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend class PairComputeFunctor; friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDSFKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDSFKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist(PairLJCutCoulDSFKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute(PairLJCutCoulDSFKokkos*, NeighListKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/cut/coul/cut/kk That style is not supported by Kokkos. */ diff --git a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h index 5bdaaf96c..732832923 100644 --- a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h +++ b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.h @@ -1,164 +1,158 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/cut/coul/long/kk,PairLJCutCoulLongKokkos) PairStyle(lj/cut/coul/long/kk/device,PairLJCutCoulLongKokkos) PairStyle(lj/cut/coul/long/kk/host,PairLJCutCoulLongKokkos) #else #ifndef LMP_PAIR_LJ_CUT_COUL_LONG_KOKKOS_H #define LMP_PAIR_LJ_CUT_COUL_LONG_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_cut_coul_long.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJCutCoulLongKokkos : public PairLJCutCoulLong { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJCutCoulLongKokkos(class LAMMPS *); ~PairLJCutCoulLongKokkos(); void compute(int, int); void settings(int, char **); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; - F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; - }; - protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; Kokkos::DualView k_params; typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCutCoulLongKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJCutCoulLongKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJCutCoulLongKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJCutCoulLongKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/cut/coul/long/kk That style is not supported by Kokkos. */ diff --git a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp index 499a82667..b636f3649 100644 --- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp +++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp @@ -1,500 +1,500 @@ /* ---------------------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ /* ---------------------------------------------------------------------- Contributing author: Ray Shan (SNL) ------------------------------------------------------------------------- */ #include #include #include #include #include "pair_lj_gromacs_coul_gromacs_kokkos.h" #include "kokkos.h" #include "atom_kokkos.h" #include "comm.h" #include "force.h" #include "neighbor.h" #include "neigh_list.h" #include "neigh_request.h" #include "update.h" #include "integrate.h" #include "respa.h" #include "math_const.h" #include "memory.h" #include "error.h" #include "atom_masks.h" using namespace LAMMPS_NS; using namespace MathConst; #define KOKKOS_CUDA_MAX_THREADS 256 #define KOKKOS_CUDA_MIN_BLOCKS 8 /* ---------------------------------------------------------------------- */ template PairLJGromacsCoulGromacsKokkos::PairLJGromacsCoulGromacsKokkos(LAMMPS *lmp):PairLJGromacsCoulGromacs(lmp) { respa_enable = 0; atomKK = (AtomKokkos *) atom; execution_space = ExecutionSpaceFromDevice::space; datamask_read = X_MASK | F_MASK | TYPE_MASK | Q_MASK | ENERGY_MASK | VIRIAL_MASK; datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK; cutsq = NULL; cut_ljsq = 0.0; cut_coulsq = 0.0; } /* ---------------------------------------------------------------------- */ template PairLJGromacsCoulGromacsKokkos::~PairLJGromacsCoulGromacsKokkos() { if (!copymode) { memory->destroy_kokkos(k_eatom,eatom); memory->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); //memory->sfree(cut_ljsq); //memory->sfree(cut_coulsq); eatom = NULL; vatom = NULL; cutsq = NULL; cut_ljsq = 0.0; cut_coulsq = 0.0; } } /* ---------------------------------------------------------------------- */ template void PairLJGromacsCoulGromacsKokkos::cleanup_copy() { // WHY needed: this prevents parent copy from deallocating any arrays allocated = 0; cutsq = NULL; cut_ljsq = 0.0; eatom = NULL; vatom = NULL; ftable = NULL; } /* ---------------------------------------------------------------------- */ template void PairLJGromacsCoulGromacsKokkos::compute(int eflag_in, int vflag_in) { eflag = eflag_in; vflag = vflag_in; if (neighflag == FULL) no_virial_fdotr_compute = 1; if (eflag || vflag) ev_setup(eflag,vflag); else evflag = vflag_fdotr = 0; atomKK->sync(execution_space,datamask_read); k_cutsq.template sync(); k_cut_ljsq.template sync(); k_cut_coulsq.template sync(); k_params.template sync(); if (eflag || vflag) atomKK->modified(execution_space,datamask_modify); else atomKK->modified(execution_space,F_MASK); x = atomKK->k_x.view(); c_x = atomKK->k_x.view(); f = atomKK->k_f.view(); q = atomKK->k_q.view(); type = atomKK->k_type.view(); nlocal = atom->nlocal; nall = atom->nlocal + atom->nghost; special_lj[0] = force->special_lj[0]; special_lj[1] = force->special_lj[1]; special_lj[2] = force->special_lj[2]; special_lj[3] = force->special_lj[3]; special_coul[0] = force->special_coul[0]; special_coul[1] = force->special_coul[1]; special_coul[2] = force->special_coul[2]; special_coul[3] = force->special_coul[3]; qqrd2e = force->qqrd2e; newton_pair = force->newton_pair; // loop over neighbors of my atoms copymode = 1; EV_FLOAT ev; if(ncoultablebits) ev = pair_compute,CoulLongTable<1> > (this,(NeighListKokkos*)list); else ev = pair_compute,CoulLongTable<0> > (this,(NeighListKokkos*)list); if (eflag) { eng_vdwl += ev.evdwl; eng_coul += ev.ecoul; } if (vflag_global) { virial[0] += ev.v[0]; virial[1] += ev.v[1]; virial[2] += ev.v[2]; virial[3] += ev.v[3]; virial[4] += ev.v[4]; virial[5] += ev.v[5]; } if (vflag_fdotr) pair_virial_fdotr_compute(this); copymode = 0; } /* ---------------------------------------------------------------------- compute LJ GROMACS pair force between atoms i and j ---------------------------------------------------------------------- */ template template KOKKOS_INLINE_FUNCTION F_FLOAT PairLJGromacsCoulGromacsKokkos:: compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const { const F_FLOAT r2inv = 1.0/rsq; const F_FLOAT r6inv = r2inv*r2inv*r2inv; F_FLOAT forcelj = r6inv * ((STACKPARAMS?m_params[itype][jtype].lj1:params(itype,jtype).lj1)*r6inv - (STACKPARAMS?m_params[itype][jtype].lj2:params(itype,jtype).lj2)); if (rsq > cut_lj_innersq) { const F_FLOAT r = sqrt(rsq); const F_FLOAT tlj = r - cut_lj_inner; const F_FLOAT fswitch = r*tlj*tlj* ((STACKPARAMS?m_params[itype][jtype].ljsw1:params(itype,jtype).ljsw1) + (STACKPARAMS?m_params[itype][jtype].ljsw2:params(itype,jtype).ljsw2)*tlj); forcelj += fswitch; } return forcelj*r2inv; } /* ---------------------------------------------------------------------- compute LJ GROMACS pair potential energy between atoms i and j ---------------------------------------------------------------------- */ template template KOKKOS_INLINE_FUNCTION F_FLOAT PairLJGromacsCoulGromacsKokkos:: compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const { const F_FLOAT r2inv = 1.0/rsq; const F_FLOAT r6inv = r2inv*r2inv*r2inv; F_FLOAT englj = r6inv * ((STACKPARAMS?m_params[itype][jtype].lj3:params(itype,jtype).lj3)*r6inv - (STACKPARAMS?m_params[itype][jtype].lj4:params(itype,jtype).lj4)); englj += (STACKPARAMS?m_params[itype][jtype].ljsw5:params(itype,jtype).ljsw5); if (rsq > cut_lj_innersq) { const F_FLOAT r = sqrt(rsq); const F_FLOAT tlj = r - cut_lj_inner; const F_FLOAT eswitch = tlj*tlj*tlj * ((STACKPARAMS?m_params[itype][jtype].ljsw3:params(itype,jtype).ljsw3) + (STACKPARAMS?m_params[itype][jtype].ljsw4:params(itype,jtype).ljsw4)*tlj); englj += eswitch; } return englj; } /* ---------------------------------------------------------------------- compute coulomb pair force between atoms i and j ---------------------------------------------------------------------- */ template template KOKKOS_INLINE_FUNCTION F_FLOAT PairLJGromacsCoulGromacsKokkos:: compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const { const F_FLOAT r2inv = 1.0/rsq; const F_FLOAT rinv = sqrt(r2inv); F_FLOAT forcecoul = qqrd2e*qtmp*q(j) *rinv; if (rsq > cut_coul_innersq) { const F_FLOAT r = 1.0/rinv; const F_FLOAT tc = r - cut_coul_inner; const F_FLOAT fcoulswitch = qqrd2e * qtmp*q(j)*r*tc*tc*(coulsw1 + coulsw2*tc); forcecoul += fcoulswitch; } return forcecoul * r2inv * factor_coul; } /* ---------------------------------------------------------------------- compute coulomb pair potential energy between atoms i and j ---------------------------------------------------------------------- */ template template KOKKOS_INLINE_FUNCTION F_FLOAT PairLJGromacsCoulGromacsKokkos:: compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const { const F_FLOAT r2inv = 1.0/rsq; const F_FLOAT rinv = sqrt(r2inv); F_FLOAT ecoul = qqrd2e * qtmp * q(j) * (rinv-coulsw5); if (rsq > cut_coul_innersq) { const F_FLOAT r = 1.0/rinv; const F_FLOAT tc = r - cut_coul_inner; const F_FLOAT ecoulswitch = tc*tc*tc * (coulsw3 + coulsw4*tc); ecoul += qqrd2e*qtmp*q(j)*ecoulswitch; } return ecoul * factor_coul; } /* ---------------------------------------------------------------------- allocate all arrays ------------------------------------------------------------------------- */ template void PairLJGromacsCoulGromacsKokkos::allocate() { PairLJGromacsCoulGromacs::allocate(); int n = atom->ntypes; memory->destroy(cutsq); memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view(); //memory->destroy(cut_ljsq); memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view(); memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view(); - k_params = Kokkos::DualView("PairLJGromacsCoulGromacs::params",n+1,n+1); + k_params = Kokkos::DualView("PairLJGromacsCoulGromacs::params",n+1,n+1); params = k_params.d_view; } template void PairLJGromacsCoulGromacsKokkos::init_tables(double cut_coul, double *cut_respa) { Pair::init_tables(cut_coul,cut_respa); typedef typename ArrayTypes::t_ffloat_1d table_type; typedef typename ArrayTypes::t_ffloat_1d host_table_type; int ntable = 1; for (int i = 0; i < ncoultablebits; i++) ntable *= 2; // Copy rtable and drtable { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); for(int i = 0; i < ntable; i++) { h_table(i) = rtable[i]; } Kokkos::deep_copy(d_table,h_table); d_rtable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); for(int i = 0; i < ntable; i++) { h_table(i) = drtable[i]; } Kokkos::deep_copy(d_table,h_table); d_drtable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); // Copy ftable and dftable for(int i = 0; i < ntable; i++) { h_table(i) = ftable[i]; } Kokkos::deep_copy(d_table,h_table); d_ftable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); for(int i = 0; i < ntable; i++) { h_table(i) = dftable[i]; } Kokkos::deep_copy(d_table,h_table); d_dftable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); // Copy ctable and dctable for(int i = 0; i < ntable; i++) { h_table(i) = ctable[i]; } Kokkos::deep_copy(d_table,h_table); d_ctable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); for(int i = 0; i < ntable; i++) { h_table(i) = dctable[i]; } Kokkos::deep_copy(d_table,h_table); d_dctable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); // Copy etable and detable for(int i = 0; i < ntable; i++) { h_table(i) = etable[i]; } Kokkos::deep_copy(d_table,h_table); d_etable = d_table; } { host_table_type h_table("HostTable",ntable); table_type d_table("DeviceTable",ntable); for(int i = 0; i < ntable; i++) { h_table(i) = detable[i]; } Kokkos::deep_copy(d_table,h_table); d_detable = d_table; } } /* ---------------------------------------------------------------------- global settings ------------------------------------------------------------------------- */ template void PairLJGromacsCoulGromacsKokkos::settings(int narg, char **arg) { if (narg > 2) error->all(FLERR,"Illegal pair_style command"); PairLJGromacsCoulGromacs::settings(narg,arg); } /* ---------------------------------------------------------------------- init specific to this pair style ------------------------------------------------------------------------- */ template void PairLJGromacsCoulGromacsKokkos::init_style() { PairLJGromacsCoulGromacs::init_style(); // error if rRESPA with inner levels if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) { int respa = 0; if (((Respa *) update->integrate)->level_inner >= 0) respa = 1; if (((Respa *) update->integrate)->level_middle >= 0) respa = 2; if (respa) error->all(FLERR,"Cannot use Kokkos pair style with rRESPA inner/middle"); } // irequest = neigh request made by parent class neighflag = lmp->kokkos->neighflag; int irequest = neighbor->nrequest - 1; neighbor->requests[irequest]-> kokkos_host = Kokkos::Impl::is_same::value && !Kokkos::Impl::is_same::value; neighbor->requests[irequest]-> kokkos_device = Kokkos::Impl::is_same::value; if (neighflag == FULL) { neighbor->requests[irequest]->full = 1; neighbor->requests[irequest]->half = 0; } else if (neighflag == HALF || neighflag == HALFTHREAD) { neighbor->requests[irequest]->full = 0; neighbor->requests[irequest]->half = 1; } else { error->all(FLERR,"Cannot use chosen neighbor list style with lj/gromacs/coul/gromacs/kk"); } } /* ---------------------------------------------------------------------- init for one type pair i,j and corresponding j,i ------------------------------------------------------------------------- */ template double PairLJGromacsCoulGromacsKokkos::init_one(int i, int j) { double cutone = PairLJGromacsCoulGromacs::init_one(i,j); double cut_ljsqm = cut_ljsq; double cut_coulsqm = cut_coulsq; k_params.h_view(i,j).lj1 = lj1[i][j]; k_params.h_view(i,j).lj2 = lj2[i][j]; k_params.h_view(i,j).lj3 = lj3[i][j]; k_params.h_view(i,j).lj4 = lj4[i][j]; k_params.h_view(i,j).ljsw1 = ljsw1[i][j]; k_params.h_view(i,j).ljsw2 = ljsw2[i][j]; k_params.h_view(i,j).ljsw3 = ljsw3[i][j]; k_params.h_view(i,j).ljsw4 = ljsw4[i][j]; k_params.h_view(i,j).ljsw5 = ljsw5[i][j]; k_params.h_view(i,j).cut_ljsq = cut_ljsqm; k_params.h_view(i,j).cut_coulsq = cut_coulsqm; k_params.h_view(j,i) = k_params.h_view(i,j); if(i(); k_cut_ljsq.h_view(i,j) = k_cut_ljsq.h_view(j,i) = cut_ljsqm; k_cut_ljsq.template modify(); k_cut_coulsq.h_view(i,j) = k_cut_coulsq.h_view(j,i) = cut_coulsqm; k_cut_coulsq.template modify(); k_params.template modify(); return cutone; } namespace LAMMPS_NS { template class PairLJGromacsCoulGromacsKokkos; #ifdef KOKKOS_HAVE_CUDA template class PairLJGromacsCoulGromacsKokkos; #endif } diff --git a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h index 8b10eb71a..bbf5c50a6 100644 --- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h +++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.h @@ -1,165 +1,167 @@ /* -*- c++ -*- ---------------------------------------------------------- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator http://lammps.sandia.gov, Sandia National Laboratories Steve Plimpton, sjplimp@sandia.gov Copyright (2003) Sandia Corporation. Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains certain rights in this software. This software is distributed under the GNU General Public License. See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ #ifdef PAIR_CLASS PairStyle(lj/gromacs/coul/gromacs/kk,PairLJGromacsCoulGromacsKokkos) PairStyle(lj/gromacs/coul/gromacs/kk/device,PairLJGromacsCoulGromacsKokkos) PairStyle(lj/gromacs/coul/gromacs/kk/host,PairLJGromacsCoulGromacsKokkos) #else #ifndef LMP_PAIR_LJ_GROMACS_COUL_GROMACS_KOKKOS_H #define LMP_PAIR_LJ_GROMACS_COUL_GROMACS_KOKKOS_H #include "pair_kokkos.h" #include "pair_lj_gromacs_coul_gromacs.h" #include "neigh_list_kokkos.h" namespace LAMMPS_NS { template class PairLJGromacsCoulGromacsKokkos : public PairLJGromacsCoulGromacs { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=1}; typedef DeviceType device_type; PairLJGromacsCoulGromacsKokkos(class LAMMPS *); ~PairLJGromacsCoulGromacsKokkos(); void compute(int, int); void settings(int, char **); void init_tables(double cut_coul, double *cut_respa); void init_style(); double init_one(int, int); - struct params_lj_coul{ - params_lj_coul(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;ljsw1=0;ljsw2=0;ljsw3=0;ljsw4=0;ljsw5=0;}; - params_lj_coul(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;ljsw1=0;ljsw2=0;ljsw3=0;ljsw4=0;ljsw5=0;}; + struct params_lj_coul_gromacs{ + KOKKOS_INLINE_FUNCTION + params_lj_coul_gromacs(){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;ljsw1=0;ljsw2=0;ljsw3=0;ljsw4=0;ljsw5=0;}; + KOKKOS_INLINE_FUNCTION + params_lj_coul_gromacs(int i){cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;ljsw1=0;ljsw2=0;ljsw3=0;ljsw4=0;ljsw5=0;}; F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset,ljsw1,ljsw2,ljsw3,ljsw4,ljsw5; }; protected: void cleanup_copy(); template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fpair(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype) const; template KOKKOS_INLINE_FUNCTION F_FLOAT compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j, const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const; - Kokkos::DualView k_params; - typename Kokkos::DualView k_params; + typename Kokkos::DualView::t_dev_const_um params; // hardwired to space for 12 atom types - params_lj_coul m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; + params_lj_coul_gromacs m_params[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cutsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_ljsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; F_FLOAT m_cut_coulsq[MAX_TYPES_STACKPARAMS+1][MAX_TYPES_STACKPARAMS+1]; typename ArrayTypes::t_x_array_randomread x; typename ArrayTypes::t_x_array c_x; typename ArrayTypes::t_f_array f; typename ArrayTypes::t_int_1d_randomread type; typename ArrayTypes::t_float_1d_randomread q; DAT::tdual_efloat_1d k_eatom; DAT::tdual_virial_array k_vatom; typename ArrayTypes::t_efloat_1d d_eatom; typename ArrayTypes::t_virial_array d_vatom; int newton_pair; typename ArrayTypes::tdual_ffloat_2d k_cutsq; typename ArrayTypes::t_ffloat_2d d_cutsq; typename ArrayTypes::tdual_ffloat_2d k_cut_ljsq; typename ArrayTypes::t_ffloat_2d d_cut_ljsq; typename ArrayTypes::tdual_ffloat_2d k_cut_coulsq; typename ArrayTypes::t_ffloat_2d d_cut_coulsq; typename ArrayTypes::t_ffloat_1d_randomread d_rtable, d_drtable, d_ftable, d_dftable, d_ctable, d_dctable, d_etable, d_detable; int neighflag; int nlocal,nall,eflag,vflag; double special_coul[4]; double special_lj[4]; double qqrd2e; void allocate(); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJGromacsCoulGromacsKokkos*, NeighListKokkos*); friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend class PairComputeFunctor >; friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute_neighlist >(PairLJGromacsCoulGromacsKokkos*,NeighListKokkos*); friend EV_FLOAT pair_compute >(PairLJGromacsCoulGromacsKokkos*, NeighListKokkos*); friend void pair_virial_fdotr_compute(PairLJGromacsCoulGromacsKokkos*); }; } #endif #endif /* ERROR/WARNING messages: E: Illegal ... command Self-explanatory. Check the input script syntax and compare to the documentation for the command. You can use -echo screen as a command-line option when running LAMMPS to see the offending line. E: Cannot use Kokkos pair style with rRESPA inner/middle Self-explanatory. E: Cannot use chosen neighbor list style with lj/gromacs/coul/gromacs/kk Self-explanatory. */