case ecutsVERLET:
set_zones_size(dd, state_local->box, &ddbox, 0, 1, ncg_moved);
- nbnxn_put_on_grid(fr->nbv, state_local->box,
+ nbnxn_put_on_grid(fr->nbv.get(), state_local->box,
0,
comm->zones.size[0].bb_x0,
comm->zones.size[0].bb_x1,
pme_load_balance(pme_lb, cr,
fp_err, fp_log, mdlog,
ir, state, pme_lb->cycles_c - cycles_prev,
- fr->ic, fr->nbv, &fr->pmedata,
+ fr->ic, fr->nbv.get(), &fr->pmedata,
step);
/* Update deprecated rlist in forcerec to stay in sync with fr->nbv */
GMX_RELEASE_ASSERT(ir->rcoulomb == ir->rvdw, "With Verlet lists and no PME rcoulomb and rvdw should be identical");
}
- Nbnxm::init_nb_verlet(mdlog, &fr->nbv, bFEP_NonBonded, ir, fr,
- cr, hardwareInfo, deviceInfo,
- mtop, box);
+ fr->nbv = Nbnxm::init_nb_verlet(mdlog, bFEP_NonBonded, ir, fr,
+ cr, hardwareInfo, deviceInfo,
+ mtop, box);
if (useGpuForBonded)
{
if (isPPrankUsingGPU)
{
- /* free nbnxn data in GPU memory */
- Nbnxm::gpu_free(fr->nbv->gpu_nbv);
+ /* Free data in GPU memory and pinned memory before destroying the GPU context */
+ fr->nbv.reset();
+
delete fr->gpuBonded;
fr->gpuBonded = nullptr;
}
return;
}
- nonbonded_verlet_t *nbv = fr->nbv;
+ nonbonded_verlet_t *nbv = fr->nbv.get();
/* GPU kernel launch overhead is already timed separately */
if (fr->cutoff_scheme != ecutsVERLET)
gmx_bool bDoForces, bUseGPU, bUseOrEmulGPU;
rvec vzero, box_diag;
float cycles_pme, cycles_wait_gpu;
- nonbonded_verlet_t *nbv = fr->nbv;
+ nonbonded_verlet_t *nbv = fr->nbv.get();
bStateChanged = ((flags & GMX_FORCE_STATECHANGED) != 0);
bNS = ((flags & GMX_FORCE_NS) != 0);
}
nbnxn_atomdata_copy_shiftvec((flags & GMX_FORCE_DYNAMICBOX) != 0,
- fr->shift_vec, nbv->nbat);
+ fr->shift_vec, nbv->nbat.get());
#if GMX_MPI
if (!thisRankHasDuty(cr, DUTY_PME))
wallcycle_sub_stop(wcycle, ewcsNBS_GRID_NONLOCAL);
}
- nbnxn_atomdata_set(nbv->nbat, nbv->nbs.get(), mdatoms, fr->cginfo);
+ nbnxn_atomdata_set(nbv->nbat.get(), nbv->nbs.get(), mdatoms, fr->cginfo);
wallcycle_stop(wcycle, ewcNS);
}
if (bNS)
{
- Nbnxm::gpu_init_atomdata(nbv->gpu_nbv, nbv->nbat);
+ Nbnxm::gpu_init_atomdata(nbv->gpu_nbv, nbv->nbat.get());
}
- Nbnxm::gpu_upload_shiftvec(nbv->gpu_nbv, nbv->nbat);
+ Nbnxm::gpu_upload_shiftvec(nbv->gpu_nbv, nbv->nbat.get());
wallcycle_sub_stop(wcycle, ewcsLAUNCH_GPU_NONBONDED);
{
nbnxn_atomdata_copy_x_to_nbat_x(nbv->nbs.get(), Nbnxm::AtomLocality::Local,
FALSE, as_rvec_array(x.unpaddedArrayRef().data()),
- nbv->nbat, wcycle);
+ nbv->nbat.get(), wcycle);
}
if (bUseGPU)
wallcycle_start(wcycle, ewcLAUNCH_GPU);
wallcycle_sub_start(wcycle, ewcsLAUNCH_GPU_NONBONDED);
- Nbnxm::gpu_copy_xq_to_gpu(nbv->gpu_nbv, nbv->nbat, Nbnxm::AtomLocality::Local, ppForceWorkload->haveGpuBondedWork);
+ Nbnxm::gpu_copy_xq_to_gpu(nbv->gpu_nbv, nbv->nbat.get(),
+ Nbnxm::AtomLocality::Local,
+ ppForceWorkload->haveGpuBondedWork);
wallcycle_sub_stop(wcycle, ewcsLAUNCH_GPU_NONBONDED);
// bonded work not split into separate local and non-local, so with DD
nbnxn_atomdata_copy_x_to_nbat_x(nbv->nbs.get(), Nbnxm::AtomLocality::NonLocal,
FALSE, as_rvec_array(x.unpaddedArrayRef().data()),
- nbv->nbat, wcycle);
+ nbv->nbat.get(), wcycle);
}
if (bUseGPU)
/* launch non-local nonbonded tasks on GPU */
wallcycle_sub_start_nocount(wcycle, ewcsLAUNCH_GPU_NONBONDED);
- Nbnxm::gpu_copy_xq_to_gpu(nbv->gpu_nbv, nbv->nbat, Nbnxm::AtomLocality::NonLocal, ppForceWorkload->haveGpuBondedWork);
+ Nbnxm::gpu_copy_xq_to_gpu(nbv->gpu_nbv, nbv->nbat.get(),
+ Nbnxm::AtomLocality::NonLocal,
+ ppForceWorkload->haveGpuBondedWork);
wallcycle_sub_stop(wcycle, ewcsLAUNCH_GPU_NONBONDED);
if (ppForceWorkload->haveGpuBondedWork)
wallcycle_sub_start_nocount(wcycle, ewcsLAUNCH_GPU_NONBONDED);
if (havePPDomainDecomposition(cr))
{
- Nbnxm::gpu_launch_cpyback(nbv->gpu_nbv, nbv->nbat,
+ Nbnxm::gpu_launch_cpyback(nbv->gpu_nbv, nbv->nbat.get(),
flags, Nbnxm::AtomLocality::NonLocal, ppForceWorkload->haveGpuBondedWork);
}
- Nbnxm::gpu_launch_cpyback(nbv->gpu_nbv, nbv->nbat,
+ Nbnxm::gpu_launch_cpyback(nbv->gpu_nbv, nbv->nbat.get(),
flags, Nbnxm::AtomLocality::Local, ppForceWorkload->haveGpuBondedWork);
wallcycle_sub_stop(wcycle, ewcsLAUNCH_GPU_NONBONDED);
{
/* This is not in a subcounter because it takes a
negligible and constant-sized amount of time */
- nbnxn_atomdata_add_nbat_fshift_to_fshift(nbv->nbat,
+ nbnxn_atomdata_add_nbat_fshift_to_fshift(nbv->nbat.get(),
fr->fshift);
}
}
bool alternateGpuWait = (!c_disableAlternatingWait && useGpuPme && bUseGPU && !DOMAINDECOMP(cr));
if (alternateGpuWait)
{
- alternatePmeNbGpuWaitReduce(fr->nbv, fr->pmedata, &force, &forceWithVirial, fr->fshift, enerd, flags, pmeFlags, ppForceWorkload->haveGpuBondedWork, wcycle);
+ alternatePmeNbGpuWaitReduce(fr->nbv.get(), fr->pmedata, &force, &forceWithVirial, fr->fshift, enerd,
+ flags, pmeFlags, ppForceWorkload->haveGpuBondedWork, wcycle);
}
if (!alternateGpuWait && useGpuPme)
if (bPMETune)
{
pme_loadbal_init(&pme_loadbal, cr, mdlog, *ir, state->box,
- *fr->ic, fr->nbv->pairlistSets().params(), fr->pmedata, use_GPU(fr->nbv),
+ *fr->ic, fr->nbv->pairlistSets().params(), fr->pmedata, use_GPU(fr->nbv.get()),
&bPMETunePrinting);
}
step_rel++;
resetHandler->resetCounters(
- step, step_rel, mdlog, fplog, cr, (use_GPU(fr->nbv) ? fr->nbv : nullptr),
+ step, step_rel, mdlog, fplog, cr, (use_GPU(fr->nbv.get()) ? fr->nbv.get() : nullptr),
nrnb, fr->pmedata, pme_loadbal, wcycle, walltime_accounting);
/* If bIMD is TRUE, the master updates the IMD energy record and sends positions to VMD client */
if (bPMETune)
{
- pme_loadbal_done(pme_loadbal, fplog, mdlog, use_GPU(fr->nbv));
+ pme_loadbal_done(pme_loadbal, fplog, mdlog, use_GPU(fr->nbv.get()));
}
done_shellfc(fplog, shellfc, step_rel);
*/
finish_run(fplog, mdlog, cr,
inputrec, nrnb, wcycle, walltime_accounting,
- fr ? fr->nbv : nullptr,
+ fr ? fr->nbv.get() : nullptr,
pmedata,
EI_DYNAMICS(inputrec->eI) && !isMultiSim(ms));
#define GMX_MDTYPES_TYPES_FORCEREC_H
#include <array>
+#include <memory>
#include <vector>
#include "gromacs/math/vectypes.h"
int cutoff_scheme = 0; /* group- or Verlet-style cutoff */
gmx_bool bNonbonded = FALSE; /* true if nonbonded calculations are *not* turned off */
- struct nonbonded_verlet_t *nbv = nullptr;
+
+ /* The Nbnxm Verlet non-bonded machinery */
+ std::unique_ptr<nonbonded_verlet_t> nbv;
/* The wall tables (if used) */
int nwall = 0;
*/
if (nbat->bUseTreeReduce)
{
- nbnxn_atomdata_add_nbat_f_to_f_treereduce(nbat, nth);
+ nbnxn_atomdata_add_nbat_f_to_f_treereduce(nbat.get(), nth);
}
else
{
- nbnxn_atomdata_add_nbat_f_to_f_stdreduce(nbat, nth);
+ nbnxn_atomdata_add_nbat_f_to_f_stdreduce(nbat.get(), nth);
}
}
#pragma omp parallel for num_threads(nth) schedule(static)
{
try
{
- nbnxn_atomdata_add_nbat_f_to_f_part(nbs.get(), nbat,
+ nbnxn_atomdata_add_nbat_f_to_f_part(nbs.get(), nbat.get(),
nbat->out,
1,
a0+((th+0)*na)/nth,
{
return;
}
- gmx_nbnxn_cuda_t *nb = nbv->gpu_nbv;
- cu_nbparam_t *nbp = nb->nbparam;
+ cu_nbparam_t *nbp = nbv->gpu_nbv->nbparam;
set_cutoff_parameters(nbp, ic, listParams);
nbp->eeltype = pick_ewald_kernel_type(ic->rcoulomb != ic->rvdw);
- init_ewald_coulomb_force_table(ic, nb->nbparam);
+ init_ewald_coulomb_force_table(ic, nbp);
}
/*! Initializes the pair list data structure. */
nbnxn_cuda_clear_e_fshift(nb);
}
-void gpu_init(gmx_nbnxn_cuda_t **p_nb,
- const gmx_device_info_t *deviceInfo,
- const interaction_const_t *ic,
- const NbnxnListParameters *listParams,
- const nbnxn_atomdata_t *nbat,
- int /*rank*/,
- gmx_bool bLocalAndNonlocal)
+gmx_nbnxn_cuda_t *
+gpu_init(const gmx_device_info_t *deviceInfo,
+ const interaction_const_t *ic,
+ const NbnxnListParameters *listParams,
+ const nbnxn_atomdata_t *nbat,
+ int /*rank*/,
+ gmx_bool bLocalAndNonlocal)
{
cudaError_t stat;
- gmx_nbnxn_cuda_t *nb;
-
- if (p_nb == nullptr)
- {
- return;
- }
+ gmx_nbnxn_cuda_t *nb;
snew(nb, 1);
snew(nb->atdat, 1);
snew(nb->nbparam, 1);
cuda_init_const(nb, ic, listParams, nbat->params());
- *p_nb = nb;
-
if (debug)
{
fprintf(debug, "Initialized CUDA data structures.\n");
}
+
+ return nb;
}
void gpu_init_pairlist(gmx_nbnxn_cuda_t *nb,
#ifndef GMX_NBNXN_GPU_DATA_MGMT_H
#define GMX_NBNXN_GPU_DATA_MGMT_H
+#include <memory>
+
#include "gromacs/gpu_utils/gpu_macros.h"
#include "gromacs/mdtypes/interaction_const.h"
/** Initializes the data structures related to GPU nonbonded calculations. */
GPU_FUNC_QUALIFIER
-void gpu_init(gmx_nbnxn_gpu_t gmx_unused **p_nb,
- const gmx_device_info_t gmx_unused *deviceInfo,
- const interaction_const_t gmx_unused *ic,
- const NbnxnListParameters gmx_unused *listParams,
- const nbnxn_atomdata_t gmx_unused *nbat,
- int gmx_unused rank,
- /* true if both local and non-local are done on GPU */
- gmx_bool gmx_unused bLocalAndNonlocal) GPU_FUNC_TERM
+gmx_nbnxn_gpu_t *
+gpu_init(const gmx_device_info_t gmx_unused *deviceInfo,
+ const interaction_const_t gmx_unused *ic,
+ const NbnxnListParameters gmx_unused *listParams,
+ const nbnxn_atomdata_t gmx_unused *nbat,
+ int gmx_unused rank,
+ /* true if both local and non-local are done on GPU */
+ gmx_bool gmx_unused bLocalAndNonlocal) GPU_FUNC_TERM_WITH_RETURN(nullptr)
/** Initializes pair-list data for GPU, called at every pair search step. */
GPU_FUNC_QUALIFIER
/** Calculates the minimum size of proximity lists to improve SM load balance
* with GPU non-bonded kernels. */
- GPU_FUNC_QUALIFIER
+ GPU_FUNC_QUALIFIER
int gpu_min_ci_balanced(gmx_nbnxn_gpu_t gmx_unused *nb) GPU_FUNC_TERM_WITH_RETURN(-1)
/** Returns if analytical Ewald GPU kernels are used. */
* Note: CUDA only.
*/
CUDA_FUNC_QUALIFIER
- rvec *gpu_get_fshift(gmx_nbnxn_gpu_t gmx_unused *nb) CUDA_FUNC_TERM_WITH_RETURN(nullptr)
+ rvec *gpu_get_fshift(gmx_nbnxn_gpu_t gmx_unused *nb) CUDA_FUNC_TERM_WITH_RETURN(nullptr)
} // namespace Nbnxm
#if GMX_GPU == GMX_GPU_OPENCL
struct gmx_nbnxn_ocl_t;
-typedef struct gmx_nbnxn_ocl_t gmx_nbnxn_gpu_t;
+using gmx_nbnxn_gpu_t = gmx_nbnxn_ocl_t;
#endif
#if GMX_GPU == GMX_GPU_CUDA
struct gmx_nbnxn_cuda_t;
-typedef struct gmx_nbnxn_cuda_t gmx_nbnxn_gpu_t;
+using gmx_nbnxn_gpu_t = gmx_nbnxn_cuda_t;
#endif
#if GMX_GPU == GMX_GPU_NONE
lowerCorner, upperCorner,
nbs->grid[0].atom_density);
- nbnxn_atomdata_t *nbat = nbv->nbat;
+ nbnxn_atomdata_t *nbat = nbv->nbat.get();
calc_cell_indices(nbs, ddZone, grid, updateGroupsCog, atomStart, atomEnd, atinfo, x, numAtomsMoved, move, nbat);
case Nbnxm::KernelType::Cpu4xN_Simd_2xNN:
nbnxn_kernel_cpu(pairlistSet,
kernelSetup(),
- nbat,
+ nbat.get(),
ic,
fr->shift_vec,
forceFlags,
case Nbnxm::KernelType::Cpu8x8x8_PlainC:
nbnxn_kernel_gpu_ref(pairlistSet.nblGpu[0],
- nbat, &ic,
+ nbat.get(), &ic,
fr->shift_vec,
forceFlags,
clearF,
int64_t outerListCreationStep_;
};
+ //! Constructs an object from its components
+ nonbonded_verlet_t(std::unique_ptr<PairlistSets> pairlistSets_,
+ std::unique_ptr<nbnxn_search> nbs,
+ std::unique_ptr<nbnxn_atomdata_t> nbat,
+ const Nbnxm::KernelSetup &kernelSetup,
+ gmx_nbnxn_gpu_t *gpu_nbv);
+
+ ~nonbonded_verlet_t();
+
//! Returns whether a GPU is use for the non-bonded calculations
bool useGpu() const
{
return kernelSetup_;
}
- //! Sets the kernel setup, TODO: make private
- void setKernelSetup(const Nbnxm::KernelSetup &kernelSetup)
- {
- kernelSetup_ = kernelSetup;
- }
-
// TODO: Make all data members private
public:
//! All data related to the pair lists
- std::unique_ptr<PairlistSets> pairlistSets_;
+ std::unique_ptr<PairlistSets> pairlistSets_;
//! Working data for constructing the pairlists
- std::unique_ptr<nbnxn_search> nbs;
+ std::unique_ptr<nbnxn_search> nbs;
//! Atom data
- nbnxn_atomdata_t *nbat;
-
+ std::unique_ptr<nbnxn_atomdata_t> nbat;
private:
//! The non-bonded setup, also affects the pairlist construction kernel
- Nbnxm::KernelSetup kernelSetup_;
+ Nbnxm::KernelSetup kernelSetup_;
public:
-
- gmx_nbnxn_gpu_t *gpu_nbv; /**< pointer to GPU nb verlet data */
+ //! GPU Nbnxm data, only used with a physical GPU (TODO: use unique_ptr)
+ gmx_nbnxn_gpu_t *gpu_nbv;
};
namespace Nbnxm
{
-/*! \brief Initializes the nbnxn module */
-void init_nb_verlet(const gmx::MDLogger &mdlog,
- nonbonded_verlet_t **nb_verlet,
- gmx_bool bFEP_NonBonded,
- const t_inputrec *ir,
- const t_forcerec *fr,
- const t_commrec *cr,
- const gmx_hw_info_t &hardwareInfo,
- const gmx_device_info_t *deviceInfo,
- const gmx_mtop_t *mtop,
- matrix box);
+/*! \brief Creates an Nbnxm object */
+std::unique_ptr<nonbonded_verlet_t>
+init_nb_verlet(const gmx::MDLogger &mdlog,
+ gmx_bool bFEP_NonBonded,
+ const t_inputrec *ir,
+ const t_forcerec *fr,
+ const t_commrec *cr,
+ const gmx_hw_info_t &hardwareInfo,
+ const gmx_device_info_t *deviceInfo,
+ const gmx_mtop_t *mtop,
+ matrix box);
} // namespace Nbnxm
#include "gromacs/utility/fatalerror.h"
#include "gromacs/utility/logger.h"
+#include "gpu_types.h"
#include "grid.h"
#include "internal.h"
return minimumIlistCount;
}
-void init_nb_verlet(const gmx::MDLogger &mdlog,
- nonbonded_verlet_t **nb_verlet,
- gmx_bool bFEP_NonBonded,
- const t_inputrec *ir,
- const t_forcerec *fr,
- const t_commrec *cr,
- const gmx_hw_info_t &hardwareInfo,
- const gmx_device_info_t *deviceInfo,
- const gmx_mtop_t *mtop,
- matrix box)
+std::unique_ptr<nonbonded_verlet_t>
+init_nb_verlet(const gmx::MDLogger &mdlog,
+ gmx_bool bFEP_NonBonded,
+ const t_inputrec *ir,
+ const t_forcerec *fr,
+ const t_commrec *cr,
+ const gmx_hw_info_t &hardwareInfo,
+ const gmx_device_info_t *deviceInfo,
+ const gmx_mtop_t *mtop,
+ matrix box)
{
- nonbonded_verlet_t *nbv = new nonbonded_verlet_t();
-
const bool emulateGpu = (getenv("GMX_EMULATE_GPU") != nullptr);
const bool useGpu = deviceInfo != nullptr;
nonbondedResource = NonbondedResource::Cpu;
}
- nbv->nbs = nullptr;
-
- nbv->setKernelSetup(pick_nbnxn_kernel(mdlog, fr->use_simd_kernels, hardwareInfo,
- nonbondedResource, ir,
- fr->bNonbonded));
+ Nbnxm::KernelSetup kernelSetup =
+ pick_nbnxn_kernel(mdlog, fr->use_simd_kernels, hardwareInfo,
+ nonbondedResource, ir,
+ fr->bNonbonded);
const bool haveMultipleDomains = (DOMAINDECOMP(cr) && cr->dd->nnodes > 1);
- NbnxnListParameters listParams(nbv->kernelSetup().kernelType,
+ NbnxnListParameters listParams(kernelSetup.kernelType,
ir->rlist,
havePPDomainDecomposition(cr));
enbnxninitcombrule = enbnxninitcombruleNONE;
}
- nbv->nbat = new nbnxn_atomdata_t(useGpu ? gmx::PinningPolicy::PinnedIfSupported : gmx::PinningPolicy::CannotBePinned);
+ std::unique_ptr<nbnxn_atomdata_t> nbat =
+ std::make_unique<nbnxn_atomdata_t>(useGpu ? gmx::PinningPolicy::PinnedIfSupported : gmx::PinningPolicy::CannotBePinned);
+
int mimimumNumEnergyGroupNonbonded = ir->opts.ngener;
if (ir->opts.ngener - ir->nwall == 1)
{
mimimumNumEnergyGroupNonbonded = 1;
}
nbnxn_atomdata_init(mdlog,
- nbv->nbat,
- nbv->kernelSetup().kernelType,
+ nbat.get(),
+ kernelSetup.kernelType,
enbnxninitcombrule,
fr->ntype, fr->nbfp,
mimimumNumEnergyGroupNonbonded,
- nbv->pairlistIsSimple() ? gmx_omp_nthreads_get(emntNonbonded) : 1);
+ (useGpu || emulateGpu) ? 1 : gmx_omp_nthreads_get(emntNonbonded));
- int minimumIlistCountForGpuBalancing = 0;
+ gmx_nbnxn_gpu_t *gpu_nbv = nullptr;
+ int minimumIlistCountForGpuBalancing = 0;
if (useGpu)
{
/* init the NxN GPU data; the last argument tells whether we'll have
* both local and non-local NB calculation on GPU */
- gpu_init(&nbv->gpu_nbv,
- deviceInfo,
- fr->ic,
- &listParams,
- nbv->nbat,
- cr->nodeid,
- haveMultipleDomains);
-
- minimumIlistCountForGpuBalancing = getMinimumIlistCountForGpuBalancing(nbv->gpu_nbv);
+ gpu_nbv = gpu_init(deviceInfo,
+ fr->ic,
+ &listParams,
+ nbat.get(),
+ cr->nodeid,
+ haveMultipleDomains);
+
+ minimumIlistCountForGpuBalancing = getMinimumIlistCountForGpuBalancing(gpu_nbv);
}
- nbv->pairlistSets_ =
+ std::unique_ptr<nonbonded_verlet_t::PairlistSets> pairlistSets =
std::make_unique<nonbonded_verlet_t::PairlistSets>(listParams,
haveMultipleDomains,
minimumIlistCountForGpuBalancing);
- nbv->nbs = std::make_unique<nbnxn_search>(ir->ePBC,
- DOMAINDECOMP(cr) ? &cr->dd->nc : nullptr,
- DOMAINDECOMP(cr) ? domdec_zones(cr->dd) : nullptr,
- bFEP_NonBonded,
- gmx_omp_nthreads_get(emntPairsearch));
-
- *nb_verlet = nbv;
+ std::unique_ptr<nbnxn_search> nbs =
+ std::make_unique<nbnxn_search>(ir->ePBC,
+ DOMAINDECOMP(cr) ? &cr->dd->nc : nullptr,
+ DOMAINDECOMP(cr) ? domdec_zones(cr->dd) : nullptr,
+ bFEP_NonBonded,
+ gmx_omp_nthreads_get(emntPairsearch));
+
+ return std::make_unique<nonbonded_verlet_t>(std::move(pairlistSets),
+ std::move(nbs),
+ std::move(nbat),
+ kernelSetup,
+ gpu_nbv);
}
} // namespace Nbnxm
+
+nonbonded_verlet_t::nonbonded_verlet_t(std::unique_ptr<PairlistSets> pairlistSets,
+ std::unique_ptr<nbnxn_search> nbs_in,
+ std::unique_ptr<nbnxn_atomdata_t> nbat_in,
+ const Nbnxm::KernelSetup &kernelSetup,
+ gmx_nbnxn_gpu_t *gpu_nbv_ptr) :
+ pairlistSets_(std::move(pairlistSets)),
+ nbs(std::move(nbs_in)),
+ nbat(std::move(nbat_in)),
+ kernelSetup_(kernelSetup),
+ gpu_nbv(gpu_nbv_ptr)
+{
+ GMX_RELEASE_ASSERT(pairlistSets_, "Need valid pairlistSets");
+ GMX_RELEASE_ASSERT(nbs, "Need valid search object");
+ GMX_RELEASE_ASSERT(nbat, "Need valid atomdata object");
+}
+
+nonbonded_verlet_t::~nonbonded_verlet_t()
+{
+ Nbnxm::gpu_free(gpu_nbv);
+}
nbp->eeltype = gpu_pick_ewald_kernel_type(ic->rcoulomb != ic->rvdw);
- init_ewald_coulomb_force_table(ic, nb->nbparam, nb->dev_rundata);
+ init_ewald_coulomb_force_table(ic, nbp, nb->dev_rundata);
}
/*! \brief Initializes the pair list data structure.
//! This function is documented in the header file
-void gpu_init(gmx_nbnxn_ocl_t **p_nb,
- const gmx_device_info_t *deviceInfo,
- const interaction_const_t *ic,
- const NbnxnListParameters *listParams,
- const nbnxn_atomdata_t *nbat,
- const int rank,
- const gmx_bool bLocalAndNonlocal)
+gmx_nbnxn_ocl_t *
+gpu_init(const gmx_device_info_t *deviceInfo,
+ const interaction_const_t *ic,
+ const NbnxnListParameters *listParams,
+ const nbnxn_atomdata_t *nbat,
+ const int rank,
+ const gmx_bool bLocalAndNonlocal)
{
gmx_nbnxn_ocl_t *nb;
cl_int cl_error;
assert(ic);
- if (p_nb == nullptr)
- {
- return;
- }
-
snew(nb, 1);
snew(nb->atdat, 1);
snew(nb->nbparam, 1);
/* clear energy and shift force outputs */
nbnxn_ocl_clear_e_fshift(nb);
- *p_nb = nb;
-
if (debug)
{
fprintf(debug, "Initialized OpenCL data structures.\n");
}
+
+ return nb;
}
/*! \brief Clears the first natoms_clear elements of the GPU nonbonded force output array.
int64_t step,
t_nrnb *nrnb)
{
- pairlistSets_->construct(iLocality, nbs.get(), nbat, excl,
+ pairlistSets_->construct(iLocality, nbs.get(), nbat.get(), excl,
kernelSetup_.kernelType,
step, nrnb);
nonbonded_verlet_t::dispatchPruneKernelCpu(const Nbnxm::InteractionLocality iLocality,
const rvec *shift_vec)
{
- pairlistSets_->dispatchPruneKernel(iLocality, nbat, shift_vec, kernelSetup_.kernelType);
+ pairlistSets_->dispatchPruneKernel(iLocality, nbat.get(), shift_vec, kernelSetup_.kernelType);
}
biod.pnpi.spb.ru / alexxy / gromacs.git / commitdiff