1 // clang-format off 2 /* -*- c++ -*- ---------------------------------------------------------- 3 LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator 4 https://www.lammps.org/, Sandia National Laboratories 5 Steve Plimpton, sjplimp@sandia.gov 6 7 Copyright (2003) Sandia Corporation. Under the terms of Contract 8 DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains 9 certain rights in this software. This software is distributed under 10 the GNU General Public License. 11 12 See the README file in the top-level LAMMPS directory. 13 ------------------------------------------------------------------------- */ 14 15 #ifndef LMP_LMPTYPE_KOKKOS_H 16 #define LMP_LMPTYPE_KOKKOS_H 17 18 #include "pointers.h" 19 #include "lmptype.h" 20 21 #include <Kokkos_Core.hpp> 22 #include <Kokkos_DualView.hpp> 23 #include <impl/Kokkos_Timer.hpp> 24 #include <Kokkos_Vectorization.hpp> 25 #include <Kokkos_ScatterView.hpp> 26 #include <Kokkos_UnorderedMap.hpp> 27 28 enum{FULL=1u,HALFTHREAD=2u,HALF=4u}; 29 30 #if defined(KOKKOS_ENABLE_CXX11) 31 #undef ISFINITE 32 #define ISFINITE(x) std::isfinite(x) 33 #endif 34 35 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP) || defined(KOKKOS_ENABLE_SYCL) 36 #define LMP_KOKKOS_GPU 37 #endif 38 39 #if defined(LMP_KOKKOS_GPU) 40 #define KOKKOS_GPU_ARG(x) x 41 #else 42 #define KOKKOS_GPU_ARG(x) 43 #endif 44 45 #define MAX_TYPES_STACKPARAMS 12 46 static constexpr LAMMPS_NS::bigint LMP_KOKKOS_AV_DELTA = 10; 47 48 namespace Kokkos { 49 static auto NoInit = [](std::string const& label) { 50 return Kokkos::view_alloc(Kokkos::WithoutInitializing, label); 51 }; 52 } 53 54 struct lmp_float3 { 55 float x,y,z; 56 KOKKOS_INLINE_FUNCTION lmp_float3lmp_float357 lmp_float3():x(0.0f),y(0.0f),z(0.0f) {} 58 59 KOKKOS_INLINE_FUNCTION 60 void operator += (const lmp_float3& tmp) { 61 x+=tmp.x; 62 y+=tmp.y; 63 z+=tmp.z; 64 } 65 KOKKOS_INLINE_FUNCTION 66 void operator += (const lmp_float3& tmp) volatile { 67 x+=tmp.x; 68 y+=tmp.y; 69 z+=tmp.z; 70 } 71 KOKKOS_INLINE_FUNCTION 72 void operator = (const lmp_float3& tmp) { 73 x=tmp.x; 74 y=tmp.y; 75 z=tmp.z; 76 } 77 KOKKOS_INLINE_FUNCTION 78 void operator = (const lmp_float3& tmp) volatile { 79 x=tmp.x; 80 y=tmp.y; 81 z=tmp.z; 82 } 83 }; 84 85 struct lmp_double3 { 86 double x,y,z; 87 KOKKOS_INLINE_FUNCTION lmp_double3lmp_double388 lmp_double3():x(0.0),y(0.0),z(0.0) {} 89 90 KOKKOS_INLINE_FUNCTION 91 void operator += (const lmp_double3& tmp) { 92 x+=tmp.x; 93 y+=tmp.y; 94 z+=tmp.z; 95 } 96 KOKKOS_INLINE_FUNCTION 97 void operator += (const lmp_double3& tmp) volatile { 98 x+=tmp.x; 99 y+=tmp.y; 100 z+=tmp.z; 101 } 102 KOKKOS_INLINE_FUNCTION 103 void operator = (const lmp_double3& tmp) { 104 x=tmp.x; 105 y=tmp.y; 106 z=tmp.z; 107 } 108 KOKKOS_INLINE_FUNCTION 109 void operator = (const lmp_double3& tmp) volatile { 110 x=tmp.x; 111 y=tmp.y; 112 z=tmp.z; 113 } 114 }; 115 116 template<class Scalar> 117 struct t_scalar3 { 118 Scalar x,y,z; 119 120 KOKKOS_FORCEINLINE_FUNCTION t_scalar3t_scalar3121 t_scalar3() { 122 x = 0; y = 0; z = 0; 123 } 124 125 KOKKOS_FORCEINLINE_FUNCTION t_scalar3t_scalar3126 t_scalar3(const t_scalar3& rhs) { 127 x = rhs.x; y = rhs.y; z = rhs.z; 128 } 129 130 KOKKOS_FORCEINLINE_FUNCTION t_scalar3t_scalar3131 t_scalar3(const Scalar& x_, const Scalar& y_, const Scalar& z_ ) { 132 x = x_; y = y_; z = z_; 133 } 134 135 KOKKOS_FORCEINLINE_FUNCTION 136 t_scalar3 operator= (const t_scalar3& rhs) { 137 x = rhs.x; y = rhs.y; z = rhs.z; 138 return *this; 139 } 140 141 KOKKOS_FORCEINLINE_FUNCTION 142 t_scalar3 operator= (const volatile t_scalar3& rhs) { 143 x = rhs.x; y = rhs.y; z = rhs.z; 144 return *this; 145 } 146 147 KOKKOS_FORCEINLINE_FUNCTION 148 t_scalar3 operator+= (const t_scalar3& rhs) { 149 x += rhs.x; y += rhs.y; z += rhs.z; 150 return *this; 151 } 152 153 KOKKOS_FORCEINLINE_FUNCTION 154 t_scalar3 operator+= (const volatile t_scalar3& rhs) volatile { 155 x += rhs.x; y += rhs.y; z += rhs.z; 156 return *this; 157 } 158 }; 159 160 template<class Scalar> 161 KOKKOS_FORCEINLINE_FUNCTION 162 t_scalar3<Scalar> operator + 163 (const t_scalar3<Scalar>& a, const t_scalar3<Scalar>& b) { 164 return t_scalar3<Scalar>(a.x+b.x,a.y+b.y,a.z+b.z); 165 } 166 167 template<class Scalar> 168 KOKKOS_FORCEINLINE_FUNCTION 169 t_scalar3<Scalar> operator * 170 (const t_scalar3<Scalar>& a, const Scalar& b) { 171 return t_scalar3<Scalar>(a.x*b,a.y*b,a.z*b); 172 } 173 174 template<class Scalar> 175 KOKKOS_FORCEINLINE_FUNCTION 176 t_scalar3<Scalar> operator * 177 (const Scalar& b, const t_scalar3<Scalar>& a) { 178 return t_scalar3<Scalar>(a.x*b,a.y*b,a.z*b); 179 } 180 181 // set LMPHostype and LMPDeviceType from Kokkos Default Types 182 typedef Kokkos::DefaultExecutionSpace LMPDeviceType; 183 typedef Kokkos::HostSpace::execution_space LMPHostType; 184 185 186 // Need to use Cuda UVM memory space for Host execution space 187 188 template<class DeviceType> 189 class KKDevice { 190 public: 191 #if defined(KOKKOS_ENABLE_CUDA) && defined(KOKKOS_ENABLE_CUDA_UVM) 192 typedef Kokkos::Device<DeviceType,LMPDeviceType::memory_space> value; 193 #else 194 typedef Kokkos::Device<DeviceType,typename DeviceType::memory_space> value; 195 #endif 196 }; 197 198 199 // set ExecutionSpace stuct with variable "space" 200 201 template<class Device> 202 struct ExecutionSpaceFromDevice; 203 204 template<> 205 struct ExecutionSpaceFromDevice<LMPHostType> { 206 static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Host; 207 }; 208 209 #ifdef KOKKOS_ENABLE_CUDA 210 template<> 211 struct ExecutionSpaceFromDevice<Kokkos::Cuda> { 212 static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Device; 213 }; 214 #elif defined(KOKKOS_ENABLE_HIP) 215 template<> 216 struct ExecutionSpaceFromDevice<Kokkos::Experimental::HIP> { 217 static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Device; 218 }; 219 #elif defined(KOKKOS_ENABLE_SYCL) 220 template<> 221 struct ExecutionSpaceFromDevice<Kokkos::Experimental::SYCL> { 222 static const LAMMPS_NS::ExecutionSpace space = LAMMPS_NS::Device; 223 }; 224 #endif 225 226 // set host pinned space 227 #if defined(KOKKOS_ENABLE_CUDA) 228 typedef Kokkos::CudaHostPinnedSpace LMPPinnedHostType; 229 #elif defined(KOKKOS_ENABLE_HIP) 230 typedef Kokkos::Experimental::HIPHostPinnedSpace LMPPinnedHostType; 231 #elif defined(KOKKOS_ENABLE_SYCL) 232 typedef Kokkos::Experimental::SYCLSharedUSMSpace LMPPinnedHostType; 233 #endif 234 235 // create simple LMPDeviceSpace typedef for non CUDA-, HIP-, or SYCL-specific 236 // behaviour 237 #if defined(KOKKOS_ENABLE_CUDA) 238 typedef Kokkos::Cuda LMPDeviceSpace; 239 #elif defined(KOKKOS_ENABLE_HIP) 240 typedef Kokkos::Experimental::HIP LMPDeviceSpace; 241 #elif defined(KOKKOS_ENABLE_SYCL) 242 typedef Kokkos::Experimental::SYCL LMPDeviceSpace; 243 #endif 244 245 246 // Determine memory traits for force array 247 // Do atomic trait when running HALFTHREAD neighbor list style 248 template<int NEIGHFLAG> 249 struct AtomicF { 250 enum {value = Kokkos::Unmanaged}; 251 }; 252 253 template<> 254 struct AtomicF<HALFTHREAD> { 255 enum {value = Kokkos::Atomic|Kokkos::Unmanaged}; 256 }; 257 258 259 // Determine memory traits for force array 260 // Do atomic trait when running HALFTHREAD neighbor list style with CUDA 261 template<int NEIGHFLAG, class DeviceType> 262 struct AtomicDup { 263 using value = Kokkos::Experimental::ScatterNonAtomic; 264 }; 265 266 #ifdef KOKKOS_ENABLE_CUDA 267 template<> 268 struct AtomicDup<HALFTHREAD,Kokkos::Cuda> { 269 using value = Kokkos::Experimental::ScatterAtomic; 270 }; 271 #elif defined(KOKKOS_ENABLE_HIP) 272 template<> 273 struct AtomicDup<HALFTHREAD,Kokkos::Experimental::HIP> { 274 using value = Kokkos::Experimental::ScatterAtomic; 275 }; 276 #elif defined(KOKKOS_ENABLE_SYCL) 277 template<> 278 struct AtomicDup<HALFTHREAD,Kokkos::Experimental::SYCL> { 279 using value = Kokkos::Experimental::ScatterAtomic; 280 }; 281 #endif 282 283 #ifdef LMP_KOKKOS_USE_ATOMICS 284 285 #ifdef KOKKOS_ENABLE_OPENMP 286 template<> 287 struct AtomicDup<HALFTHREAD,Kokkos::OpenMP> { 288 using value = Kokkos::Experimental::ScatterAtomic; 289 }; 290 #endif 291 292 #ifdef KOKKOS_ENABLE_THREADS 293 template<> 294 struct AtomicDup<HALFTHREAD,Kokkos::Threads> { 295 using value = Kokkos::Experimental::ScatterAtomic; 296 }; 297 #endif 298 299 #endif 300 301 302 // Determine duplication traits for force array 303 // Use duplication when running threaded and not using atomics 304 template<int NEIGHFLAG, class DeviceType> 305 struct NeedDup { 306 using value = Kokkos::Experimental::ScatterNonDuplicated; 307 }; 308 309 #ifndef LMP_KOKKOS_USE_ATOMICS 310 311 #ifdef KOKKOS_ENABLE_OPENMP 312 template<> 313 struct NeedDup<HALFTHREAD,Kokkos::OpenMP> { 314 using value = Kokkos::Experimental::ScatterDuplicated; 315 }; 316 #endif 317 318 #ifdef KOKKOS_ENABLE_THREADS 319 template<> 320 struct NeedDup<HALFTHREAD,Kokkos::Threads> { 321 using value = Kokkos::Experimental::ScatterDuplicated; 322 }; 323 #endif 324 325 #endif 326 327 template<typename value, typename T1, typename T2> 328 class ScatterViewHelper {}; 329 330 template<typename T1, typename T2> 331 class ScatterViewHelper<Kokkos::Experimental::ScatterDuplicated,T1,T2> { 332 public: 333 KOKKOS_INLINE_FUNCTION 334 static T1 get(const T1 &dup, const T2 & /*nondup*/) { 335 return dup; 336 } 337 }; 338 339 template<typename T1, typename T2> 340 class ScatterViewHelper<Kokkos::Experimental::ScatterNonDuplicated,T1,T2> { 341 public: 342 KOKKOS_INLINE_FUNCTION 343 static T2 get(const T1 & /*dup*/, const T2 &nondup) { 344 return nondup; 345 } 346 }; 347 348 349 // define precision 350 // handle global precision, force, energy, positions, kspace separately 351 352 #ifndef PRECISION 353 #define PRECISION 2 354 #endif 355 #if PRECISION==1 356 typedef float LMP_FLOAT; 357 #else 358 typedef double LMP_FLOAT; 359 #endif 360 361 #ifndef PREC_FORCE 362 #define PREC_FORCE PRECISION 363 #endif 364 365 #if PREC_FORCE==1 366 typedef float F_FLOAT; 367 #else 368 typedef double F_FLOAT; 369 #endif 370 371 #ifndef PREC_ENERGY 372 #define PREC_ENERGY PRECISION 373 #endif 374 375 #if PREC_ENERGY==1 376 typedef float E_FLOAT; 377 #else 378 typedef double E_FLOAT; 379 #endif 380 381 struct s_EV_FLOAT { 382 E_FLOAT evdwl; 383 E_FLOAT ecoul; 384 E_FLOAT v[6]; 385 KOKKOS_INLINE_FUNCTION 386 s_EV_FLOAT() { 387 evdwl = 0; 388 ecoul = 0; 389 v[0] = 0; v[1] = 0; v[2] = 0; 390 v[3] = 0; v[4] = 0; v[5] = 0; 391 } 392 393 KOKKOS_INLINE_FUNCTION 394 void operator+=(const s_EV_FLOAT &rhs) { 395 evdwl += rhs.evdwl; 396 ecoul += rhs.ecoul; 397 v[0] += rhs.v[0]; 398 v[1] += rhs.v[1]; 399 v[2] += rhs.v[2]; 400 v[3] += rhs.v[3]; 401 v[4] += rhs.v[4]; 402 v[5] += rhs.v[5]; 403 } 404 405 KOKKOS_INLINE_FUNCTION 406 void operator+=(const volatile s_EV_FLOAT &rhs) volatile { 407 evdwl += rhs.evdwl; 408 ecoul += rhs.ecoul; 409 v[0] += rhs.v[0]; 410 v[1] += rhs.v[1]; 411 v[2] += rhs.v[2]; 412 v[3] += rhs.v[3]; 413 v[4] += rhs.v[4]; 414 v[5] += rhs.v[5]; 415 } 416 }; 417 typedef struct s_EV_FLOAT EV_FLOAT; 418 419 struct s_EV_FLOAT_REAX { 420 E_FLOAT evdwl; 421 E_FLOAT ecoul; 422 E_FLOAT v[6]; 423 E_FLOAT ereax[9]; 424 KOKKOS_INLINE_FUNCTION 425 s_EV_FLOAT_REAX() { 426 evdwl = 0; 427 ecoul = 0; 428 v[0] = 0; v[1] = 0; v[2] = 0; 429 v[3] = 0; v[4] = 0; v[5] = 0; 430 ereax[0] = 0; ereax[1] = 0; ereax[2] = 0; 431 ereax[3] = 0; ereax[4] = 0; ereax[5] = 0; 432 ereax[6] = 0; ereax[7] = 0; ereax[8] = 0; 433 } 434 435 KOKKOS_INLINE_FUNCTION 436 void operator+=(const s_EV_FLOAT_REAX &rhs) { 437 evdwl += rhs.evdwl; 438 ecoul += rhs.ecoul; 439 v[0] += rhs.v[0]; 440 v[1] += rhs.v[1]; 441 v[2] += rhs.v[2]; 442 v[3] += rhs.v[3]; 443 v[4] += rhs.v[4]; 444 v[5] += rhs.v[5]; 445 ereax[0] += rhs.ereax[0]; 446 ereax[1] += rhs.ereax[1]; 447 ereax[2] += rhs.ereax[2]; 448 ereax[3] += rhs.ereax[3]; 449 ereax[4] += rhs.ereax[4]; 450 ereax[5] += rhs.ereax[5]; 451 ereax[6] += rhs.ereax[6]; 452 ereax[7] += rhs.ereax[7]; 453 ereax[8] += rhs.ereax[8]; 454 } 455 456 KOKKOS_INLINE_FUNCTION 457 void operator+=(const volatile s_EV_FLOAT_REAX &rhs) volatile { 458 evdwl += rhs.evdwl; 459 ecoul += rhs.ecoul; 460 v[0] += rhs.v[0]; 461 v[1] += rhs.v[1]; 462 v[2] += rhs.v[2]; 463 v[3] += rhs.v[3]; 464 v[4] += rhs.v[4]; 465 v[5] += rhs.v[5]; 466 ereax[0] += rhs.ereax[0]; 467 ereax[1] += rhs.ereax[1]; 468 ereax[2] += rhs.ereax[2]; 469 ereax[3] += rhs.ereax[3]; 470 ereax[4] += rhs.ereax[4]; 471 ereax[5] += rhs.ereax[5]; 472 ereax[6] += rhs.ereax[6]; 473 ereax[7] += rhs.ereax[7]; 474 ereax[8] += rhs.ereax[8]; 475 } 476 }; 477 typedef struct s_EV_FLOAT_REAX EV_FLOAT_REAX; 478 479 struct s_FEV_FLOAT { 480 F_FLOAT f[3]; 481 E_FLOAT evdwl; 482 E_FLOAT ecoul; 483 E_FLOAT v[6]; 484 KOKKOS_INLINE_FUNCTION 485 s_FEV_FLOAT() { 486 f[0] = 0; f[1] = 0; f[2] = 0; 487 evdwl = 0; 488 ecoul = 0; 489 v[0] = 0; v[1] = 0; v[2] = 0; 490 v[3] = 0; v[4] = 0; v[5] = 0; 491 } 492 493 KOKKOS_INLINE_FUNCTION 494 void operator+=(const s_FEV_FLOAT &rhs) { 495 f[0] += rhs.f[0]; 496 f[1] += rhs.f[1]; 497 f[2] += rhs.f[2]; 498 evdwl += rhs.evdwl; 499 ecoul += rhs.ecoul; 500 v[0] += rhs.v[0]; 501 v[1] += rhs.v[1]; 502 v[2] += rhs.v[2]; 503 v[3] += rhs.v[3]; 504 v[4] += rhs.v[4]; 505 v[5] += rhs.v[5]; 506 } 507 508 KOKKOS_INLINE_FUNCTION 509 void operator+=(const volatile s_FEV_FLOAT &rhs) volatile { 510 f[0] += rhs.f[0]; 511 f[1] += rhs.f[1]; 512 f[2] += rhs.f[2]; 513 evdwl += rhs.evdwl; 514 ecoul += rhs.ecoul; 515 v[0] += rhs.v[0]; 516 v[1] += rhs.v[1]; 517 v[2] += rhs.v[2]; 518 v[3] += rhs.v[3]; 519 v[4] += rhs.v[4]; 520 v[5] += rhs.v[5]; 521 } 522 }; 523 typedef struct s_FEV_FLOAT FEV_FLOAT; 524 525 #ifndef PREC_POS 526 #define PREC_POS PRECISION 527 #endif 528 529 #if PREC_POS==1 530 typedef float X_FLOAT; 531 #else 532 typedef double X_FLOAT; 533 #endif 534 535 #ifndef PREC_VELOCITIES 536 #define PREC_VELOCITIES PRECISION 537 #endif 538 539 #if PREC_VELOCITIES==1 540 typedef float V_FLOAT; 541 #else 542 typedef double V_FLOAT; 543 #endif 544 545 #if PREC_KSPACE==1 546 typedef float K_FLOAT; 547 #else 548 typedef double K_FLOAT; 549 #endif 550 551 typedef int T_INT; 552 553 // ------------------------------------------------------------------------ 554 555 // LAMMPS types 556 557 typedef Kokkos::UnorderedMap<LAMMPS_NS::tagint,int,LMPDeviceType> hash_type; 558 typedef hash_type::HostMirror host_hash_type; 559 560 struct dual_hash_type { 561 hash_type d_view; 562 host_hash_type h_view; 563 564 template<class DeviceType> 565 KOKKOS_INLINE_FUNCTION 566 std::enable_if_t<(std::is_same<DeviceType,LMPDeviceType>::value || Kokkos::SpaceAccessibility<LMPDeviceType::memory_space,LMPHostType::memory_space>::accessible),hash_type&> view() {return d_view;} 567 568 template<class DeviceType> 569 KOKKOS_INLINE_FUNCTION 570 std::enable_if_t<!(std::is_same<DeviceType,LMPDeviceType>::value || Kokkos::SpaceAccessibility<LMPDeviceType::memory_space,LMPHostType::memory_space>::accessible),host_hash_type&> view() {return h_view;} 571 572 }; 573 574 template <class DeviceType> 575 struct ArrayTypes; 576 577 template <> 578 struct ArrayTypes<LMPDeviceType> { 579 580 // scalar types 581 582 typedef Kokkos:: 583 DualView<int, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_scalar; 584 typedef tdual_int_scalar::t_dev t_int_scalar; 585 typedef tdual_int_scalar::t_dev_const t_int_scalar_const; 586 typedef tdual_int_scalar::t_dev_um t_int_scalar_um; 587 typedef tdual_int_scalar::t_dev_const_um t_int_scalar_const_um; 588 589 typedef Kokkos:: 590 DualView<LMP_FLOAT, LMPDeviceType::array_layout, LMPDeviceType> 591 tdual_float_scalar; 592 typedef tdual_float_scalar::t_dev t_float_scalar; 593 typedef tdual_float_scalar::t_dev_const t_float_scalar_const; 594 typedef tdual_float_scalar::t_dev_um t_float_scalar_um; 595 typedef tdual_float_scalar::t_dev_const_um t_float_scalar_const_um; 596 597 // generic array types 598 599 typedef Kokkos:: 600 DualView<int*, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_1d; 601 typedef tdual_int_1d::t_dev t_int_1d; 602 typedef tdual_int_1d::t_dev_const t_int_1d_const; 603 typedef tdual_int_1d::t_dev_um t_int_1d_um; 604 typedef tdual_int_1d::t_dev_const_um t_int_1d_const_um; 605 typedef tdual_int_1d::t_dev_const_randomread t_int_1d_randomread; 606 607 typedef Kokkos:: 608 DualView<int*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_int_1d_3; 609 typedef tdual_int_1d_3::t_dev t_int_1d_3; 610 typedef tdual_int_1d_3::t_dev_const t_int_1d_3_const; 611 typedef tdual_int_1d_3::t_dev_um t_int_1d_3_um; 612 typedef tdual_int_1d_3::t_dev_const_um t_int_1d_3_const_um; 613 typedef tdual_int_1d_3::t_dev_const_randomread t_int_1d_3_randomread; 614 615 typedef Kokkos:: 616 DualView<int**, Kokkos::LayoutRight, LMPDeviceType> tdual_int_2d; 617 typedef tdual_int_2d::t_dev t_int_2d; 618 typedef tdual_int_2d::t_dev_const t_int_2d_const; 619 typedef tdual_int_2d::t_dev_um t_int_2d_um; 620 typedef tdual_int_2d::t_dev_const_um t_int_2d_const_um; 621 typedef tdual_int_2d::t_dev_const_randomread t_int_2d_randomread; 622 623 typedef Kokkos:: 624 DualView<int**, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_2d_dl; 625 typedef tdual_int_2d_dl::t_dev t_int_2d_dl; 626 typedef tdual_int_2d_dl::t_dev_const t_int_2d_const_dl; 627 typedef tdual_int_2d_dl::t_dev_um t_int_2d_um_dl; 628 typedef tdual_int_2d_dl::t_dev_const_um t_int_2d_const_um_dl; 629 typedef tdual_int_2d_dl::t_dev_const_randomread t_int_2d_randomread_dl; 630 631 typedef Kokkos:: 632 DualView<LAMMPS_NS::tagint*, LMPDeviceType::array_layout, LMPDeviceType> 633 tdual_tagint_1d; 634 typedef tdual_tagint_1d::t_dev t_tagint_1d; 635 typedef tdual_tagint_1d::t_dev_const t_tagint_1d_const; 636 typedef tdual_tagint_1d::t_dev_um t_tagint_1d_um; 637 typedef tdual_tagint_1d::t_dev_const_um t_tagint_1d_const_um; 638 typedef tdual_tagint_1d::t_dev_const_randomread t_tagint_1d_randomread; 639 640 typedef Kokkos:: 641 DualView<LAMMPS_NS::tagint**, Kokkos::LayoutRight, LMPDeviceType> 642 tdual_tagint_2d; 643 typedef tdual_tagint_2d::t_dev t_tagint_2d; 644 typedef tdual_tagint_2d::t_dev_const t_tagint_2d_const; 645 typedef tdual_tagint_2d::t_dev_um t_tagint_2d_um; 646 typedef tdual_tagint_2d::t_dev_const_um t_tagint_2d_const_um; 647 typedef tdual_tagint_2d::t_dev_const_randomread t_tagint_2d_randomread; 648 649 typedef Kokkos:: 650 DualView<LAMMPS_NS::imageint*, LMPDeviceType::array_layout, LMPDeviceType> 651 tdual_imageint_1d; 652 typedef tdual_imageint_1d::t_dev t_imageint_1d; 653 typedef tdual_imageint_1d::t_dev_const t_imageint_1d_const; 654 typedef tdual_imageint_1d::t_dev_um t_imageint_1d_um; 655 typedef tdual_imageint_1d::t_dev_const_um t_imageint_1d_const_um; 656 typedef tdual_imageint_1d::t_dev_const_randomread t_imageint_1d_randomread; 657 658 typedef Kokkos:: 659 DualView<double*, Kokkos::LayoutRight, LMPDeviceType> tdual_double_1d; 660 typedef tdual_double_1d::t_dev t_double_1d; 661 typedef tdual_double_1d::t_dev_const t_double_1d_const; 662 typedef tdual_double_1d::t_dev_um t_double_1d_um; 663 typedef tdual_double_1d::t_dev_const_um t_double_1d_const_um; 664 typedef tdual_double_1d::t_dev_const_randomread t_double_1d_randomread; 665 666 typedef Kokkos:: 667 DualView<double**, Kokkos::LayoutRight, LMPDeviceType> tdual_double_2d; 668 typedef tdual_double_2d::t_dev t_double_2d; 669 typedef tdual_double_2d::t_dev_const t_double_2d_const; 670 typedef tdual_double_2d::t_dev_um t_double_2d_um; 671 typedef tdual_double_2d::t_dev_const_um t_double_2d_const_um; 672 typedef tdual_double_2d::t_dev_const_randomread t_double_2d_randomread; 673 674 // 1d float array n 675 676 typedef Kokkos::DualView<LMP_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_float_1d; 677 typedef tdual_float_1d::t_dev t_float_1d; 678 typedef tdual_float_1d::t_dev_const t_float_1d_const; 679 typedef tdual_float_1d::t_dev_um t_float_1d_um; 680 typedef tdual_float_1d::t_dev_const_um t_float_1d_const_um; 681 typedef tdual_float_1d::t_dev_const_randomread t_float_1d_randomread; 682 683 //2d float array n 684 typedef Kokkos::DualView<LMP_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_float_2d; 685 typedef tdual_float_2d::t_dev t_float_2d; 686 typedef tdual_float_2d::t_dev_const t_float_2d_const; 687 typedef tdual_float_2d::t_dev_um t_float_2d_um; 688 typedef tdual_float_2d::t_dev_const_um t_float_2d_const_um; 689 typedef tdual_float_2d::t_dev_const_randomread t_float_2d_randomread; 690 691 //Position Types 692 //1d X_FLOAT array n 693 typedef Kokkos::DualView<X_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_xfloat_1d; 694 typedef tdual_xfloat_1d::t_dev t_xfloat_1d; 695 typedef tdual_xfloat_1d::t_dev_const t_xfloat_1d_const; 696 typedef tdual_xfloat_1d::t_dev_um t_xfloat_1d_um; 697 typedef tdual_xfloat_1d::t_dev_const_um t_xfloat_1d_const_um; 698 typedef tdual_xfloat_1d::t_dev_const_randomread t_xfloat_1d_randomread; 699 700 //2d X_FLOAT array n*m 701 typedef Kokkos::DualView<X_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_xfloat_2d; 702 typedef tdual_xfloat_2d::t_dev t_xfloat_2d; 703 typedef tdual_xfloat_2d::t_dev_const t_xfloat_2d_const; 704 typedef tdual_xfloat_2d::t_dev_um t_xfloat_2d_um; 705 typedef tdual_xfloat_2d::t_dev_const_um t_xfloat_2d_const_um; 706 typedef tdual_xfloat_2d::t_dev_const_randomread t_xfloat_2d_randomread; 707 708 //2d X_FLOAT array n*4 709 #ifdef LMP_KOKKOS_NO_LEGACY 710 typedef Kokkos::DualView<X_FLOAT*[3], Kokkos::LayoutLeft, LMPDeviceType> tdual_x_array; 711 #else 712 typedef Kokkos::DualView<X_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_x_array; 713 #endif 714 typedef tdual_x_array::t_dev t_x_array; 715 typedef tdual_x_array::t_dev_const t_x_array_const; 716 typedef tdual_x_array::t_dev_um t_x_array_um; 717 typedef tdual_x_array::t_dev_const_um t_x_array_const_um; 718 typedef tdual_x_array::t_dev_const_randomread t_x_array_randomread; 719 720 //Velocity Types 721 //1d V_FLOAT array n 722 typedef Kokkos::DualView<V_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_vfloat_1d; 723 typedef tdual_vfloat_1d::t_dev t_vfloat_1d; 724 typedef tdual_vfloat_1d::t_dev_const t_vfloat_1d_const; 725 typedef tdual_vfloat_1d::t_dev_um t_vfloat_1d_um; 726 typedef tdual_vfloat_1d::t_dev_const_um t_vfloat_1d_const_um; 727 typedef tdual_vfloat_1d::t_dev_const_randomread t_vfloat_1d_randomread; 728 729 //2d V_FLOAT array n*m 730 typedef Kokkos::DualView<V_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_vfloat_2d; 731 typedef tdual_vfloat_2d::t_dev t_vfloat_2d; 732 typedef tdual_vfloat_2d::t_dev_const t_vfloat_2d_const; 733 typedef tdual_vfloat_2d::t_dev_um t_vfloat_2d_um; 734 typedef tdual_vfloat_2d::t_dev_const_um t_vfloat_2d_const_um; 735 typedef tdual_vfloat_2d::t_dev_const_randomread t_vfloat_2d_randomread; 736 737 //2d V_FLOAT array n*3 738 typedef Kokkos::DualView<V_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_v_array; 739 //typedef Kokkos::DualView<V_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> tdual_v_array; 740 typedef tdual_v_array::t_dev t_v_array; 741 typedef tdual_v_array::t_dev_const t_v_array_const; 742 typedef tdual_v_array::t_dev_um t_v_array_um; 743 typedef tdual_v_array::t_dev_const_um t_v_array_const_um; 744 typedef tdual_v_array::t_dev_const_randomread t_v_array_randomread; 745 746 //Force Types 747 //1d F_FLOAT array n 748 749 typedef Kokkos::DualView<F_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_ffloat_1d; 750 typedef tdual_ffloat_1d::t_dev t_ffloat_1d; 751 typedef tdual_ffloat_1d::t_dev_const t_ffloat_1d_const; 752 typedef tdual_ffloat_1d::t_dev_um t_ffloat_1d_um; 753 typedef tdual_ffloat_1d::t_dev_const_um t_ffloat_1d_const_um; 754 typedef tdual_ffloat_1d::t_dev_const_randomread t_ffloat_1d_randomread; 755 756 //2d F_FLOAT array n*m 757 758 typedef Kokkos::DualView<F_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_ffloat_2d; 759 typedef tdual_ffloat_2d::t_dev t_ffloat_2d; 760 typedef tdual_ffloat_2d::t_dev_const t_ffloat_2d_const; 761 typedef tdual_ffloat_2d::t_dev_um t_ffloat_2d_um; 762 typedef tdual_ffloat_2d::t_dev_const_um t_ffloat_2d_const_um; 763 typedef tdual_ffloat_2d::t_dev_const_randomread t_ffloat_2d_randomread; 764 765 //2d F_FLOAT array n*m, device layout 766 767 typedef Kokkos::DualView<F_FLOAT**, LMPDeviceType::array_layout, LMPDeviceType> tdual_ffloat_2d_dl; 768 typedef tdual_ffloat_2d_dl::t_dev t_ffloat_2d_dl; 769 typedef tdual_ffloat_2d_dl::t_dev_const t_ffloat_2d_const_dl; 770 typedef tdual_ffloat_2d_dl::t_dev_um t_ffloat_2d_um_dl; 771 typedef tdual_ffloat_2d_dl::t_dev_const_um t_ffloat_2d_const_um_dl; 772 typedef tdual_ffloat_2d_dl::t_dev_const_randomread t_ffloat_2d_randomread_dl; 773 774 //2d F_FLOAT array n*3 775 776 typedef Kokkos::DualView<F_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_f_array; 777 //typedef Kokkos::DualView<F_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> tdual_f_array; 778 typedef tdual_f_array::t_dev t_f_array; 779 typedef tdual_f_array::t_dev_const t_f_array_const; 780 typedef tdual_f_array::t_dev_um t_f_array_um; 781 typedef tdual_f_array::t_dev_const_um t_f_array_const_um; 782 typedef tdual_f_array::t_dev_const_randomread t_f_array_randomread; 783 784 //2d F_FLOAT array n*6 (for virial) 785 786 typedef Kokkos::DualView<F_FLOAT*[6], Kokkos::LayoutRight, LMPDeviceType> tdual_virial_array; 787 typedef tdual_virial_array::t_dev t_virial_array; 788 typedef tdual_virial_array::t_dev_const t_virial_array_const; 789 typedef tdual_virial_array::t_dev_um t_virial_array_um; 790 typedef tdual_virial_array::t_dev_const_um t_virial_array_const_um; 791 typedef tdual_virial_array::t_dev_const_randomread t_virial_array_randomread; 792 793 // Spin Types 794 795 //3d SP_FLOAT array n*4 796 #ifdef LMP_KOKKOS_NO_LEGACY 797 typedef Kokkos::DualView<X_FLOAT*[4], Kokkos::LayoutLeft, LMPDeviceType> tdual_float_1d_4; 798 #else 799 typedef Kokkos::DualView<X_FLOAT*[4], Kokkos::LayoutRight, LMPDeviceType> tdual_float_1d_4; 800 #endif 801 typedef tdual_float_1d_4::t_dev t_sp_array; 802 typedef tdual_float_1d_4::t_dev_const t_sp_array_const; 803 typedef tdual_float_1d_4::t_dev_um t_sp_array_um; 804 typedef tdual_float_1d_4::t_dev_const_um t_sp_array_const_um; 805 typedef tdual_float_1d_4::t_dev_const_randomread t_sp_array_randomread; 806 807 //3d FM_FLOAT array n*3 808 809 typedef tdual_f_array::t_dev t_fm_array; 810 typedef tdual_f_array::t_dev_const t_fm_array_const; 811 typedef tdual_f_array::t_dev_um t_fm_array_um; 812 typedef tdual_f_array::t_dev_const_um t_fm_array_const_um; 813 typedef tdual_f_array::t_dev_const_randomread t_fm_array_randomread; 814 815 //3d FML_FLOAT array n*3 816 817 typedef tdual_f_array::t_dev t_fm_long_array; 818 typedef tdual_f_array::t_dev_const t_fm_long_array_const; 819 typedef tdual_f_array::t_dev_um t_fm_long_array_um; 820 typedef tdual_f_array::t_dev_const_um t_fm_long_array_const_um; 821 typedef tdual_f_array::t_dev_const_randomread t_fm_long_array_randomread; 822 823 //Energy Types 824 //1d E_FLOAT array n 825 826 typedef Kokkos::DualView<E_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_efloat_1d; 827 typedef tdual_efloat_1d::t_dev t_efloat_1d; 828 typedef tdual_efloat_1d::t_dev_const t_efloat_1d_const; 829 typedef tdual_efloat_1d::t_dev_um t_efloat_1d_um; 830 typedef tdual_efloat_1d::t_dev_const_um t_efloat_1d_const_um; 831 typedef tdual_efloat_1d::t_dev_const_randomread t_efloat_1d_randomread; 832 833 //2d E_FLOAT array n*m 834 835 typedef Kokkos::DualView<E_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_efloat_2d; 836 typedef tdual_efloat_2d::t_dev t_efloat_2d; 837 typedef tdual_efloat_2d::t_dev_const t_efloat_2d_const; 838 typedef tdual_efloat_2d::t_dev_um t_efloat_2d_um; 839 typedef tdual_efloat_2d::t_dev_const_um t_efloat_2d_const_um; 840 typedef tdual_efloat_2d::t_dev_const_randomread t_efloat_2d_randomread; 841 842 //2d E_FLOAT array n*3 843 844 typedef Kokkos::DualView<E_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_e_array; 845 typedef tdual_e_array::t_dev t_e_array; 846 typedef tdual_e_array::t_dev_const t_e_array_const; 847 typedef tdual_e_array::t_dev_um t_e_array_um; 848 typedef tdual_e_array::t_dev_const_um t_e_array_const_um; 849 typedef tdual_e_array::t_dev_const_randomread t_e_array_randomread; 850 851 //Neighbor Types 852 853 typedef Kokkos::DualView<int**, LMPDeviceType::array_layout, LMPDeviceType> tdual_neighbors_2d; 854 typedef tdual_neighbors_2d::t_dev t_neighbors_2d; 855 typedef tdual_neighbors_2d::t_dev_const t_neighbors_2d_const; 856 typedef tdual_neighbors_2d::t_dev_um t_neighbors_2d_um; 857 typedef tdual_neighbors_2d::t_dev_const_um t_neighbors_2d_const_um; 858 typedef tdual_neighbors_2d::t_dev_const_randomread t_neighbors_2d_randomread; 859 860 }; 861 862 #ifdef LMP_KOKKOS_GPU 863 template <> 864 struct ArrayTypes<LMPHostType> { 865 866 //Scalar Types 867 868 typedef Kokkos::DualView<int, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_scalar; 869 typedef tdual_int_scalar::t_host t_int_scalar; 870 typedef tdual_int_scalar::t_host_const t_int_scalar_const; 871 typedef tdual_int_scalar::t_host_um t_int_scalar_um; 872 typedef tdual_int_scalar::t_host_const_um t_int_scalar_const_um; 873 874 typedef Kokkos::DualView<LMP_FLOAT, LMPDeviceType::array_layout, LMPDeviceType> tdual_float_scalar; 875 typedef tdual_float_scalar::t_host t_float_scalar; 876 typedef tdual_float_scalar::t_host_const t_float_scalar_const; 877 typedef tdual_float_scalar::t_host_um t_float_scalar_um; 878 typedef tdual_float_scalar::t_host_const_um t_float_scalar_const_um; 879 880 //Generic ArrayTypes 881 typedef Kokkos::DualView<int*, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_1d; 882 typedef tdual_int_1d::t_host t_int_1d; 883 typedef tdual_int_1d::t_host_const t_int_1d_const; 884 typedef tdual_int_1d::t_host_um t_int_1d_um; 885 typedef tdual_int_1d::t_host_const_um t_int_1d_const_um; 886 typedef tdual_int_1d::t_host_const_randomread t_int_1d_randomread; 887 888 typedef Kokkos::DualView<int*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_int_1d_3; 889 typedef tdual_int_1d_3::t_host t_int_1d_3; 890 typedef tdual_int_1d_3::t_host_const t_int_1d_3_const; 891 typedef tdual_int_1d_3::t_host_um t_int_1d_3_um; 892 typedef tdual_int_1d_3::t_host_const_um t_int_1d_3_const_um; 893 typedef tdual_int_1d_3::t_host_const_randomread t_int_1d_3_randomread; 894 895 typedef Kokkos::DualView<int**, Kokkos::LayoutRight, LMPDeviceType> tdual_int_2d; 896 typedef tdual_int_2d::t_host t_int_2d; 897 typedef tdual_int_2d::t_host_const t_int_2d_const; 898 typedef tdual_int_2d::t_host_um t_int_2d_um; 899 typedef tdual_int_2d::t_host_const_um t_int_2d_const_um; 900 typedef tdual_int_2d::t_host_const_randomread t_int_2d_randomread; 901 902 typedef Kokkos::DualView<int**, LMPDeviceType::array_layout, LMPDeviceType> tdual_int_2d_dl; 903 typedef tdual_int_2d_dl::t_host t_int_2d_dl; 904 typedef tdual_int_2d_dl::t_host_const t_int_2d_const_dl; 905 typedef tdual_int_2d_dl::t_host_um t_int_2d_um_dl; 906 typedef tdual_int_2d_dl::t_host_const_um t_int_2d_const_um_dl; 907 typedef tdual_int_2d_dl::t_host_const_randomread t_int_2d_randomread_dl; 908 909 typedef Kokkos::DualView<LAMMPS_NS::tagint*, LMPDeviceType::array_layout, LMPDeviceType> tdual_tagint_1d; 910 typedef tdual_tagint_1d::t_host t_tagint_1d; 911 typedef tdual_tagint_1d::t_host_const t_tagint_1d_const; 912 typedef tdual_tagint_1d::t_host_um t_tagint_1d_um; 913 typedef tdual_tagint_1d::t_host_const_um t_tagint_1d_const_um; 914 typedef tdual_tagint_1d::t_host_const_randomread t_tagint_1d_randomread; 915 916 typedef Kokkos:: 917 DualView<LAMMPS_NS::tagint**, Kokkos::LayoutRight, LMPDeviceType> 918 tdual_tagint_2d; 919 typedef tdual_tagint_2d::t_host t_tagint_2d; 920 typedef tdual_tagint_2d::t_host_const t_tagint_2d_const; 921 typedef tdual_tagint_2d::t_host_um t_tagint_2d_um; 922 typedef tdual_tagint_2d::t_host_const_um t_tagint_2d_const_um; 923 typedef tdual_tagint_2d::t_host_const_randomread t_tagint_2d_randomread; 924 925 typedef Kokkos:: 926 DualView<LAMMPS_NS::imageint*, LMPDeviceType::array_layout, LMPDeviceType> 927 tdual_imageint_1d; 928 typedef tdual_imageint_1d::t_host t_imageint_1d; 929 typedef tdual_imageint_1d::t_host_const t_imageint_1d_const; 930 typedef tdual_imageint_1d::t_host_um t_imageint_1d_um; 931 typedef tdual_imageint_1d::t_host_const_um t_imageint_1d_const_um; 932 typedef tdual_imageint_1d::t_host_const_randomread t_imageint_1d_randomread; 933 934 typedef Kokkos:: 935 DualView<double*, Kokkos::LayoutRight, LMPDeviceType> tdual_double_1d; 936 typedef tdual_double_1d::t_host t_double_1d; 937 typedef tdual_double_1d::t_host_const t_double_1d_const; 938 typedef tdual_double_1d::t_host_um t_double_1d_um; 939 typedef tdual_double_1d::t_host_const_um t_double_1d_const_um; 940 typedef tdual_double_1d::t_host_const_randomread t_double_1d_randomread; 941 942 typedef Kokkos:: 943 DualView<double**, Kokkos::LayoutRight, LMPDeviceType> tdual_double_2d; 944 typedef tdual_double_2d::t_host t_double_2d; 945 typedef tdual_double_2d::t_host_const t_double_2d_const; 946 typedef tdual_double_2d::t_host_um t_double_2d_um; 947 typedef tdual_double_2d::t_host_const_um t_double_2d_const_um; 948 typedef tdual_double_2d::t_host_const_randomread t_double_2d_randomread; 949 950 //1d float array n 951 typedef Kokkos::DualView<LMP_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_float_1d; 952 typedef tdual_float_1d::t_host t_float_1d; 953 typedef tdual_float_1d::t_host_const t_float_1d_const; 954 typedef tdual_float_1d::t_host_um t_float_1d_um; 955 typedef tdual_float_1d::t_host_const_um t_float_1d_const_um; 956 typedef tdual_float_1d::t_host_const_randomread t_float_1d_randomread; 957 958 //2d float array n 959 typedef Kokkos::DualView<LMP_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_float_2d; 960 typedef tdual_float_2d::t_host t_float_2d; 961 typedef tdual_float_2d::t_host_const t_float_2d_const; 962 typedef tdual_float_2d::t_host_um t_float_2d_um; 963 typedef tdual_float_2d::t_host_const_um t_float_2d_const_um; 964 typedef tdual_float_2d::t_host_const_randomread t_float_2d_randomread; 965 966 //Position Types 967 //1d X_FLOAT array n 968 typedef Kokkos::DualView<X_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_xfloat_1d; 969 typedef tdual_xfloat_1d::t_host t_xfloat_1d; 970 typedef tdual_xfloat_1d::t_host_const t_xfloat_1d_const; 971 typedef tdual_xfloat_1d::t_host_um t_xfloat_1d_um; 972 typedef tdual_xfloat_1d::t_host_const_um t_xfloat_1d_const_um; 973 typedef tdual_xfloat_1d::t_host_const_randomread t_xfloat_1d_randomread; 974 975 //2d X_FLOAT array n*m 976 typedef Kokkos::DualView<X_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_xfloat_2d; 977 typedef tdual_xfloat_2d::t_host t_xfloat_2d; 978 typedef tdual_xfloat_2d::t_host_const t_xfloat_2d_const; 979 typedef tdual_xfloat_2d::t_host_um t_xfloat_2d_um; 980 typedef tdual_xfloat_2d::t_host_const_um t_xfloat_2d_const_um; 981 typedef tdual_xfloat_2d::t_host_const_randomread t_xfloat_2d_randomread; 982 983 //2d X_FLOAT array n*3 984 typedef Kokkos::DualView<X_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_x_array; 985 typedef tdual_x_array::t_host t_x_array; 986 typedef tdual_x_array::t_host_const t_x_array_const; 987 typedef tdual_x_array::t_host_um t_x_array_um; 988 typedef tdual_x_array::t_host_const_um t_x_array_const_um; 989 typedef tdual_x_array::t_host_const_randomread t_x_array_randomread; 990 991 //Velocity Types 992 //1d V_FLOAT array n 993 typedef Kokkos::DualView<V_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_vfloat_1d; 994 typedef tdual_vfloat_1d::t_host t_vfloat_1d; 995 typedef tdual_vfloat_1d::t_host_const t_vfloat_1d_const; 996 typedef tdual_vfloat_1d::t_host_um t_vfloat_1d_um; 997 typedef tdual_vfloat_1d::t_host_const_um t_vfloat_1d_const_um; 998 typedef tdual_vfloat_1d::t_host_const_randomread t_vfloat_1d_randomread; 999 1000 //2d V_FLOAT array n*m 1001 typedef Kokkos::DualView<V_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_vfloat_2d; 1002 typedef tdual_vfloat_2d::t_host t_vfloat_2d; 1003 typedef tdual_vfloat_2d::t_host_const t_vfloat_2d_const; 1004 typedef tdual_vfloat_2d::t_host_um t_vfloat_2d_um; 1005 typedef tdual_vfloat_2d::t_host_const_um t_vfloat_2d_const_um; 1006 typedef tdual_vfloat_2d::t_host_const_randomread t_vfloat_2d_randomread; 1007 1008 //2d V_FLOAT array n*3 1009 typedef Kokkos::DualView<V_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_v_array; 1010 //typedef Kokkos::DualView<V_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> tdual_v_array; 1011 typedef tdual_v_array::t_host t_v_array; 1012 typedef tdual_v_array::t_host_const t_v_array_const; 1013 typedef tdual_v_array::t_host_um t_v_array_um; 1014 typedef tdual_v_array::t_host_const_um t_v_array_const_um; 1015 typedef tdual_v_array::t_host_const_randomread t_v_array_randomread; 1016 1017 //Force Types 1018 //1d F_FLOAT array n 1019 typedef Kokkos::DualView<F_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_ffloat_1d; 1020 typedef tdual_ffloat_1d::t_host t_ffloat_1d; 1021 typedef tdual_ffloat_1d::t_host_const t_ffloat_1d_const; 1022 typedef tdual_ffloat_1d::t_host_um t_ffloat_1d_um; 1023 typedef tdual_ffloat_1d::t_host_const_um t_ffloat_1d_const_um; 1024 typedef tdual_ffloat_1d::t_host_const_randomread t_ffloat_1d_randomread; 1025 1026 //2d F_FLOAT array n*m 1027 typedef Kokkos::DualView<F_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_ffloat_2d; 1028 typedef tdual_ffloat_2d::t_host t_ffloat_2d; 1029 typedef tdual_ffloat_2d::t_host_const t_ffloat_2d_const; 1030 typedef tdual_ffloat_2d::t_host_um t_ffloat_2d_um; 1031 typedef tdual_ffloat_2d::t_host_const_um t_ffloat_2d_const_um; 1032 typedef tdual_ffloat_2d::t_host_const_randomread t_ffloat_2d_randomread; 1033 1034 //2d F_FLOAT array n*m, device layout 1035 typedef Kokkos::DualView<F_FLOAT**, LMPDeviceType::array_layout, LMPDeviceType> tdual_ffloat_2d_dl; 1036 typedef tdual_ffloat_2d_dl::t_host t_ffloat_2d_dl; 1037 typedef tdual_ffloat_2d_dl::t_host_const t_ffloat_2d_const_dl; 1038 typedef tdual_ffloat_2d_dl::t_host_um t_ffloat_2d_um_dl; 1039 typedef tdual_ffloat_2d_dl::t_host_const_um t_ffloat_2d_const_um_dl; 1040 typedef tdual_ffloat_2d_dl::t_host_const_randomread t_ffloat_2d_randomread_dl; 1041 1042 //2d F_FLOAT array n*3 1043 typedef Kokkos::DualView<F_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_f_array; 1044 //typedef Kokkos::DualView<F_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> tdual_f_array; 1045 typedef tdual_f_array::t_host t_f_array; 1046 typedef tdual_f_array::t_host_const t_f_array_const; 1047 typedef tdual_f_array::t_host_um t_f_array_um; 1048 typedef tdual_f_array::t_host_const_um t_f_array_const_um; 1049 typedef tdual_f_array::t_host_const_randomread t_f_array_randomread; 1050 1051 //2d F_FLOAT array n*6 (for virial) 1052 typedef Kokkos::DualView<F_FLOAT*[6], Kokkos::LayoutRight, LMPDeviceType> tdual_virial_array; 1053 typedef tdual_virial_array::t_host t_virial_array; 1054 typedef tdual_virial_array::t_host_const t_virial_array_const; 1055 typedef tdual_virial_array::t_host_um t_virial_array_um; 1056 typedef tdual_virial_array::t_host_const_um t_virial_array_const_um; 1057 typedef tdual_virial_array::t_host_const_randomread t_virial_array_randomread; 1058 1059 // Spin types 1060 1061 //2d X_FLOAT array n*4 1062 #ifdef LMP_KOKKOS_NO_LEGACY 1063 typedef Kokkos::DualView<X_FLOAT*[4], Kokkos::LayoutLeft, LMPDeviceType> tdual_float_1d_4; 1064 #else 1065 typedef Kokkos::DualView<X_FLOAT*[4], Kokkos::LayoutRight, LMPDeviceType> tdual_float_1d_4; 1066 #endif 1067 typedef tdual_float_1d_4::t_host t_sp_array; 1068 typedef tdual_float_1d_4::t_host_const t_sp_array_const; 1069 typedef tdual_float_1d_4::t_host_um t_sp_array_um; 1070 typedef tdual_float_1d_4::t_host_const_um t_sp_array_const_um; 1071 typedef tdual_float_1d_4::t_host_const_randomread t_sp_array_randomread; 1072 1073 //2d F_FLOAT array n*3 1074 typedef tdual_f_array::t_host t_fm_array; 1075 typedef tdual_f_array::t_host_const t_fm_array_const; 1076 typedef tdual_f_array::t_host_um t_fm_array_um; 1077 typedef tdual_f_array::t_host_const_um t_fm_array_const_um; 1078 typedef tdual_f_array::t_host_const_randomread t_fm_array_randomread; 1079 1080 //2d F_FLOAT array n*3 1081 typedef tdual_f_array::t_host t_fm_long_array; 1082 typedef tdual_f_array::t_host_const t_fm_long_array_const; 1083 typedef tdual_f_array::t_host_um t_fm_long_array_um; 1084 typedef tdual_f_array::t_host_const_um t_fm_long_array_const_um; 1085 typedef tdual_f_array::t_host_const_randomread t_fm_long_array_randomread; 1086 1087 1088 //Energy Types 1089 //1d E_FLOAT array n 1090 typedef Kokkos::DualView<E_FLOAT*, LMPDeviceType::array_layout, LMPDeviceType> tdual_efloat_1d; 1091 typedef tdual_efloat_1d::t_host t_efloat_1d; 1092 typedef tdual_efloat_1d::t_host_const t_efloat_1d_const; 1093 typedef tdual_efloat_1d::t_host_um t_efloat_1d_um; 1094 typedef tdual_efloat_1d::t_host_const_um t_efloat_1d_const_um; 1095 typedef tdual_efloat_1d::t_host_const_randomread t_efloat_1d_randomread; 1096 1097 //2d E_FLOAT array n*m 1098 typedef Kokkos::DualView<E_FLOAT**, Kokkos::LayoutRight, LMPDeviceType> tdual_efloat_2d; 1099 typedef tdual_efloat_2d::t_host t_efloat_2d; 1100 typedef tdual_efloat_2d::t_host_const t_efloat_2d_const; 1101 typedef tdual_efloat_2d::t_host_um t_efloat_2d_um; 1102 typedef tdual_efloat_2d::t_host_const_um t_efloat_2d_const_um; 1103 typedef tdual_efloat_2d::t_host_const_randomread t_efloat_2d_randomread; 1104 1105 //2d E_FLOAT array n*3 1106 typedef Kokkos::DualView<E_FLOAT*[3], Kokkos::LayoutRight, LMPDeviceType> tdual_e_array; 1107 typedef tdual_e_array::t_host t_e_array; 1108 typedef tdual_e_array::t_host_const t_e_array_const; 1109 typedef tdual_e_array::t_host_um t_e_array_um; 1110 typedef tdual_e_array::t_host_const_um t_e_array_const_um; 1111 typedef tdual_e_array::t_host_const_randomread t_e_array_randomread; 1112 1113 //Neighbor Types 1114 typedef Kokkos::DualView<int**, LMPDeviceType::array_layout, LMPDeviceType> tdual_neighbors_2d; 1115 typedef tdual_neighbors_2d::t_host t_neighbors_2d; 1116 typedef tdual_neighbors_2d::t_host_const t_neighbors_2d_const; 1117 typedef tdual_neighbors_2d::t_host_um t_neighbors_2d_um; 1118 typedef tdual_neighbors_2d::t_host_const_um t_neighbors_2d_const_um; 1119 typedef tdual_neighbors_2d::t_host_const_randomread t_neighbors_2d_randomread; 1120 1121 }; 1122 #endif 1123 //default LAMMPS Types 1124 typedef struct ArrayTypes<LMPDeviceType> DAT; 1125 typedef struct ArrayTypes<LMPHostType> HAT; 1126 1127 template<class DeviceType, class BufferView, class DualView> 1128 void buffer_view(BufferView &buf, DualView &view, 1129 const size_t n0, 1130 const size_t n1) { 1131 1132 buf = BufferView(view.template view<DeviceType>().data(),n0,n1); 1133 1134 } 1135 1136 template<class DeviceType> 1137 struct MemsetZeroFunctor { 1138 typedef DeviceType execution_space ; 1139 void* ptr; 1140 KOKKOS_INLINE_FUNCTION void operator()(const int i) const { 1141 ((int*)ptr)[i] = 0; 1142 } 1143 }; 1144 1145 template<class ViewType> 1146 void memset_kokkos (ViewType &view) { 1147 static MemsetZeroFunctor<typename ViewType::execution_space> f; 1148 f.ptr = view.data(); 1149 #ifndef KOKKOS_USING_DEPRECATED_VIEW 1150 Kokkos::parallel_for(view.span()*sizeof(typename ViewType::value_type)/4, f); 1151 #else 1152 Kokkos::parallel_for(view.span()*sizeof(typename ViewType::value_type)/4, f); 1153 #endif 1154 ViewType::execution_space().fence(); 1155 } 1156 1157 struct params_lj_coul { 1158 KOKKOS_INLINE_FUNCTION 1159 params_lj_coul() {cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; 1160 KOKKOS_INLINE_FUNCTION 1161 params_lj_coul(int /*i*/) {cut_ljsq=0;cut_coulsq=0;lj1=0;lj2=0;lj3=0;lj4=0;offset=0;}; 1162 F_FLOAT cut_ljsq,cut_coulsq,lj1,lj2,lj3,lj4,offset; 1163 }; 1164 1165 // Pair SNAP 1166 1167 #define SNAP_KOKKOS_REAL double 1168 #define SNAP_KOKKOS_HOST_VECLEN 1 1169 1170 #ifdef LMP_KOKKOS_GPU 1171 #define SNAP_KOKKOS_DEVICE_VECLEN 32 1172 #else 1173 #define SNAP_KOKKOS_DEVICE_VECLEN 1 1174 #endif 1175 1176 1177 // intentional: SNAreal/complex gets reused beyond SNAP 1178 typedef double SNAreal; 1179 1180 //typedef struct { SNAreal re, im; } SNAcomplex; 1181 template <typename real_type_> 1182 struct alignas(2*sizeof(real_type_)) SNAComplex 1183 { 1184 using real_type = real_type_; 1185 using complex = SNAComplex<real_type>; 1186 real_type re,im; 1187 1188 KOKKOS_FORCEINLINE_FUNCTION SNAComplex() 1189 : re(static_cast<real_type>(0.)), im(static_cast<real_type>(0.)) { ; } 1190 1191 KOKKOS_FORCEINLINE_FUNCTION SNAComplex(real_type re) 1192 : re(re), im(static_cast<real_type>(0.)) { ; } 1193 1194 KOKKOS_FORCEINLINE_FUNCTION SNAComplex(real_type re, real_type im) 1195 : re(re), im(im) { ; } 1196 1197 KOKKOS_FORCEINLINE_FUNCTION SNAComplex(const SNAComplex& other) 1198 : re(other.re), im(other.im) { ; } 1199 1200 KOKKOS_FORCEINLINE_FUNCTION SNAComplex& operator=(const SNAComplex& other) { 1201 re = other.re; im = other.im; 1202 return *this; 1203 } 1204 1205 KOKKOS_FORCEINLINE_FUNCTION SNAComplex(SNAComplex&& other) 1206 : re(other.re), im(other.im) { ; } 1207 1208 KOKKOS_FORCEINLINE_FUNCTION SNAComplex& operator=(SNAComplex&& other) { 1209 re = other.re; im = other.im; 1210 return *this; 1211 } 1212 1213 KOKKOS_FORCEINLINE_FUNCTION SNAComplex operator+(SNAComplex const& other) { 1214 return SNAComplex(re + other.re, im + other.im); 1215 } 1216 1217 KOKKOS_FORCEINLINE_FUNCTION SNAComplex& operator+=(SNAComplex const& other) { 1218 re += other.re; im += other.im; 1219 return *this; 1220 } 1221 1222 KOKKOS_INLINE_FUNCTION 1223 static constexpr complex zero() { return complex(static_cast<real_type>(0.), static_cast<real_type>(0.)); } 1224 1225 KOKKOS_INLINE_FUNCTION 1226 static constexpr complex one() { return complex(static_cast<real_type>(1.), static_cast<real_type>(0.)); } 1227 1228 KOKKOS_INLINE_FUNCTION 1229 const complex conj() { return complex(re, -im); } 1230 1231 }; 1232 1233 template <typename real_type> 1234 KOKKOS_FORCEINLINE_FUNCTION SNAComplex<real_type> operator*(const real_type& r, const SNAComplex<real_type>& self) { 1235 return SNAComplex<real_type>(r*self.re, r*self.im); 1236 } 1237 1238 typedef SNAComplex<SNAreal> SNAcomplex; 1239 1240 #if defined(KOKKOS_ENABLE_CXX11) 1241 #undef ISFINITE 1242 #define ISFINITE(x) std::isfinite(x) 1243 #endif 1244 1245 #define LAMMPS_LAMBDA KOKKOS_LAMBDA 1246 1247 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP) 1248 #define LAMMPS_DEVICE_FUNCTION __device__ 1249 #else 1250 #define LAMMPS_DEVICE_FUNCTION 1251 #endif 1252 1253 #ifdef LMP_KOKKOS_GPU 1254 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) || defined(__SYCL_DEVICE_ONLY__) 1255 #define LMP_KK_DEVICE_COMPILE 1256 #endif 1257 #endif 1258 1259 #endif 1260