biod.pnpi.spb.ru
/
alexxy
/
gromacs.git
/ blobdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
|
commitdiff
|
tree
raw
|
inline
| side by side
SYCL: Use acc.bind(cgh) instead of cgh.require(acc)
[alexxy/gromacs.git]
/
src
/
gromacs
/
nbnxm
/
sycl
/
nbnxm_sycl_kernel.cpp
diff --git
a/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
b/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
index ea5321645d826439d21e1e0689b432c5f04caedc..a7b2c6cde6b080ac4946855703037d4209aab454 100644
(file)
--- a/
src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
+++ b/
src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
@@
-583,34
+583,34
@@
auto nbnxmKernel(cl::sycl::handler& cgh,
{
static constexpr EnergyFunctionProperties<elecType, vdwType> props;
{
static constexpr EnergyFunctionProperties<elecType, vdwType> props;
-
cgh.require(a_xq
);
-
cgh.require(a_f
);
-
cgh.require(a_shiftVec
);
-
cgh.require(a_fShift
);
+
a_xq.bind(cgh
);
+
a_f.bind(cgh
);
+
a_shiftVec.bind(cgh
);
+
a_fShift.bind(cgh
);
if constexpr (doCalcEnergies)
{
if constexpr (doCalcEnergies)
{
-
cgh.require(a_energyElec
);
-
cgh.require(a_energyVdw
);
+
a_energyElec.bind(cgh
);
+
a_energyVdw.bind(cgh
);
}
}
-
cgh.require(a_plistCJ4
);
-
cgh.require(a_plistSci
);
-
cgh.require(a_plistExcl
);
+
a_plistCJ4.bind(cgh
);
+
a_plistSci.bind(cgh
);
+
a_plistExcl.bind(cgh
);
if constexpr (!props.vdwComb)
{
if constexpr (!props.vdwComb)
{
-
cgh.require(a_atomTypes
);
-
cgh.require(a_nbfp
);
+
a_atomTypes.bind(cgh
);
+
a_nbfp.bind(cgh
);
}
else
{
}
else
{
-
cgh.require(a_ljComb
);
+
a_ljComb.bind(cgh
);
}
if constexpr (props.vdwEwald)
{
}
if constexpr (props.vdwEwald)
{
-
cgh.require(a_nbfpComb
);
+
a_nbfpComb.bind(cgh
);
}
if constexpr (props.elecEwaldTab)
{
}
if constexpr (props.elecEwaldTab)
{
-
cgh.require(a_coulombTab
);
+
a_coulombTab.bind(cgh
);
}
// shmem buffer for i x+q pre-loading
}
// shmem buffer for i x+q pre-loading