//! Number of work-items in a work-group
constexpr static int sc_workGroupSize = 256;
+//! \brief Function returning the SETTLE kernel lambda.
template<bool updateVelocities, bool computeVirial>
auto settleKernel(cl::sycl::handler& cgh,
const int numSettles,
DeviceAccessor<Float3, mode::read_write> a_xp,
float invdt,
OptionalAccessor<Float3, mode::read_write, updateVelocities> a_v,
- OptionalAccessor<float, mode_atomic, computeVirial> a_virialScaled,
+ OptionalAccessor<float, mode::read_write, computeVirial> a_virialScaled,
PbcAiuc pbcAiuc)
{
- cgh.require(a_settles);
- cgh.require(a_x);
- cgh.require(a_xp);
+ a_settles.bind(cgh);
+ a_x.bind(cgh);
+ a_xp.bind(cgh);
if constexpr (updateVelocities)
{
- cgh.require(a_v);
+ a_v.bind(cgh);
}
if constexpr (computeVirial)
{
- cgh.require(a_virialScaled);
+ a_virialScaled.bind(cgh);
}
// shmem buffer for i x+q pre-loading
// First 6 threads in the block add the 6 components of virial to the global memory address
if (threadIdx < 6)
{
- atomicFetchAdd(a_virialScaled, threadIdx, sm_threadVirial[threadIdx * blockSize]);
+ atomicFetchAdd(a_virialScaled[threadIdx], sm_threadVirial[threadIdx * blockSize]);
}
}
};
template<bool updateVelocities, bool computeVirial>
class SettleKernelName;
+//! \brief SETTLE SYCL kernel launch code.
template<bool updateVelocities, bool computeVirial, class... Args>
static cl::sycl::event launchSettleKernel(const DeviceStream& deviceStream, int numSettles, Args&&... args)
{