Page Menu
Home
c4science
Search
Configure Global Search
Log In
Files
F67306158
kokkos_type.h
No One
Temporary
Actions
Download File
Edit File
Delete File
View Transforms
Subscribe
Mute Notifications
Award Token
Subscribers
None
File Metadata
Details
File Info
Storage
Attached
Created
Fri, Jun 21, 16:30
Size
33 KB
Mime Type
text/x-c++
Expires
Sun, Jun 23, 16:30 (2 d)
Engine
blob
Format
Raw Data
Handle
18329583
Attached To
rLAMMPS lammps
kokkos_type.h
View Options
/* -*- 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 <Kokkos_Core.hpp>
#include <Kokkos_DualView.hpp>
#include <impl/Kokkos_Timer.hpp>
#include <Kokkos_Vectorization.hpp>
#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
<
class
Device
>
struct
ExecutionSpaceFromDevice
;
template
<>
struct
ExecutionSpaceFromDevice
<
LMPHostType
>
{
static
const
LAMMPS_NS
::
ExecutionSpace
space
=
LAMMPS_NS
::
Host
;
};
#ifdef KOKKOS_HAVE_CUDA
template
<>
struct
ExecutionSpaceFromDevice
<
Kokkos
::
Cuda
>
{
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
<
class
DeviceType
>
struct
ArrayTypes
;
template
<>
struct
ArrayTypes
<
LMPDeviceType
>
{
// scalar types
typedef
Kokkos
::
DualView
<
int
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMP_FLOAT
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
int
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
int
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
tagint
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
tagint
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
imageint
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
double
*
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
double
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
LMP_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMP_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
X_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
X_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
X_FLOAT
*
[
3
],
Kokkos
::
LayoutLeft
,
LMPDeviceType
>
tdual_x_array
;
#else
typedef
Kokkos
::
DualView
<
X_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
V_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
V_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
V_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_v_array
;
//typedef Kokkos::DualView<V_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> 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
<
F_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
F_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
F_FLOAT
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
F_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_f_array
;
//typedef Kokkos::DualView<F_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> 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
<
F_FLOAT
*
[
6
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
E_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
E_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
E_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
FFT_SCALAR
*
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
FFT_SCALAR
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_2d
;
typedef
tdual_FFT_SCALAR_2d
::
t_dev
t_FFT_SCALAR_2d
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
**
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_2d_3
;
typedef
tdual_FFT_SCALAR_2d_3
::
t_dev
t_FFT_SCALAR_2d_3
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
***
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_3d
;
typedef
tdual_FFT_SCALAR_3d
::
t_dev
t_FFT_SCALAR_3d
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
*
[
2
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMPHostType
>
{
//Scalar Types
typedef
Kokkos
::
DualView
<
int
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMP_FLOAT
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
int
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
int
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
tagint
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
tagint
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
LAMMPS_NS
::
imageint
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
double
*
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
double
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
LMP_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMP_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
X_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
X_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
X_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
V_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
V_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
V_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_v_array
;
//typedef Kokkos::DualView<V_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> 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
<
F_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
F_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
F_FLOAT
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
F_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_f_array
;
//typedef Kokkos::DualView<F_FLOAT*[3], LMPDeviceType::array_layout, LMPDeviceType> 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
<
F_FLOAT
*
[
6
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
E_FLOAT
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
E_FLOAT
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
E_FLOAT
*
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
**
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
FFT_SCALAR
*
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
FFT_SCALAR
**
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_2d
;
typedef
tdual_FFT_SCALAR_2d
::
t_host
t_FFT_SCALAR_2d
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
**
[
3
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_2d_3
;
typedef
tdual_FFT_SCALAR_2d_3
::
t_host
t_FFT_SCALAR_2d_3
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
***
,
Kokkos
::
LayoutRight
,
LMPDeviceType
>
tdual_FFT_SCALAR_3d
;
typedef
tdual_FFT_SCALAR_3d
::
t_host
t_FFT_SCALAR_3d
;
typedef
Kokkos
::
DualView
<
FFT_SCALAR
*
[
2
],
Kokkos
::
LayoutRight
,
LMPDeviceType
>
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
<
int
*
,
LMPDeviceType
::
array_layout
,
LMPDeviceType
>
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
<
LMPDeviceType
>
DAT
;
typedef
struct
ArrayTypes
<
LMPHostType
>
HAT
;
template
<
class
DeviceType
,
class
BufferView
,
class
DualView
>
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
<
DeviceType
>
().
ptr_on_device
(),
n0
,
n1
,
n2
,
n3
,
n4
,
n5
,
n6
,
n7
);
}
template
<
class
DeviceType
>
struct
MemsetZeroFunctor
{
typedef
DeviceType
execution_space
;
void
*
ptr
;
KOKKOS_INLINE_FUNCTION
void
operator
()(
const
int
i
)
const
{
((
int
*
)
ptr
)[
i
]
=
0
;
}
};
template
<
class
ViewType
>
void
memset_kokkos
(
ViewType
&
view
)
{
static
MemsetZeroFunctor
<
typename
ViewType
::
execution_space
>
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
Event Timeline
Log In to Comment