* \author Anca Hamuraru <anca@streamcomputing.eu>
* \author Dimitrios Karkoulis <dimitris.karkoulis@gmail.com>
* \author Teemu Virolainen <teemu@streamcomputing.eu>
+ * \author Szilárd Páll <pall.szilard@gmail.com>
*/
#include "gmxpre.h"
#include <stdlib.h>
#include <string.h>
-#include "gromacs/gmxlib/gpu_utils/gpu_utils.h"
-#include "gromacs/gmxlib/ocl_tools/oclutils.h"
-#include "gromacs/legacyheaders/gmx_detect_hardware.h"
-#include "gromacs/legacyheaders/typedefs.h"
-#include "gromacs/legacyheaders/types/enums.h"
-#include "gromacs/legacyheaders/types/force_flags.h"
-#include "gromacs/legacyheaders/types/interaction_const.h"
+#include "gromacs/gpu_utils/gpu_utils.h"
+#include "gromacs/gpu_utils/oclutils.h"
+#include "gromacs/hardware/detecthardware.h"
+#include "gromacs/hardware/gpu_hw_info.h"
+#include "gromacs/math/vectypes.h"
+#include "gromacs/mdlib/force_flags.h"
#include "gromacs/mdlib/nb_verlet.h"
#include "gromacs/mdlib/nbnxn_consts.h"
#include "gromacs/mdlib/nbnxn_gpu.h"
#include "gromacs/mdlib/nbnxn_gpu_data_mgmt.h"
#include "gromacs/mdlib/nbnxn_gpu_jit_support.h"
+#include "gromacs/mdtypes/interaction_const.h"
+#include "gromacs/mdtypes/md_enums.h"
#include "gromacs/pbcutil/ishift.h"
#include "gromacs/timing/gpu_timing.h"
#include "gromacs/utility/cstringutil.h"
#include "gromacs/utility/fatalerror.h"
+#include "gromacs/utility/gmxassert.h"
+#include "gromacs/utility/real.h"
#include "gromacs/utility/smalloc.h"
+#include "nbnxn_ocl_internal.h"
#include "nbnxn_ocl_types.h"
-
/*! \brief This parameter should be determined heuristically from the
* kernel execution times
*
*/
static unsigned int gpu_min_ci_balanced_factor = 50;
-/*! \brief Helper function for warning output
+
+/*! \brief Returns true if LJ combination rules are used in the non-bonded kernels.
*
- * We should actually be using md_print_warn in md_logging.c,
- * but we can't include mpi.h in OpenCL code.
- */
-static void md_print_warn(FILE *fplog,
- const char *fmt, ...)
+ * Full doc in nbnxn_ocl_internal.h */
+bool useLjCombRule(int vdwType)
{
- va_list ap;
-
- if (fplog != NULL)
- {
- /* We should only print to stderr on the master node,
- * in most cases fplog is only set on the master node, so this works.
- */
- va_start(ap, fmt);
- fprintf(stderr, "\n");
- vfprintf(stderr, fmt, ap);
- fprintf(stderr, "\n");
- va_end(ap);
-
- va_start(ap, fmt);
- fprintf(fplog, "\n");
- vfprintf(fplog, fmt, ap);
- fprintf(fplog, "\n");
- va_end(ap);
- }
+ return (vdwType == evdwOclCUTCOMBGEOM ||
+ vdwType == evdwOclCUTCOMBLB);
}
/*! \brief Free device buffers
bool bAsync = true,
cl_event *copy_event = NULL)
{
- cl_int cl_error;
-
if (d_dest == NULL || req_size < 0)
{
return;
than the current requested size */
if (req_size > *curr_alloc_size)
{
+ cl_int gmx_unused cl_error;
+
/* only free if the array has already been initialized */
if (*curr_alloc_size >= 0)
{
* If called with an already allocated table, it just re-uploads the
* table.
*/
-static void init_ewald_coulomb_force_table(const interaction_const_t *ic,
- cl_nbparam_t *nbp,
- const gmx_device_info_t *dev_info)
+static void init_ewald_coulomb_force_table(const interaction_const_t *ic,
+ cl_nbparam_t *nbp,
+ const gmx_device_runtime_data_t *runData)
{
cl_mem coul_tab;
array_format.image_channel_data_type = CL_FLOAT;
array_format.image_channel_order = CL_R;
- coul_tab = clCreateImage2D(dev_info->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+ coul_tab = clCreateImage2D(runData->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
&array_format, tabsize, 1, 0, ftmp, &cl_error);
*/
- coul_tab = clCreateBuffer(dev_info->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, ic->tabq_size*sizeof(cl_float), ic->tabq_coul_F, &cl_error);
+ coul_tab = clCreateBuffer(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, ic->tabq_size*sizeof(cl_float), ic->tabq_coul_F, &cl_error);
assert(cl_error == CL_SUCCESS);
// TODO: handle errors, check clCreateBuffer flags
/*! \brief Initializes the atomdata structure first time, it only gets filled at
pair-search.
*/
-static void init_atomdata_first(cl_atomdata_t *ad, int ntypes, gmx_device_info_t *dev_info)
+static void init_atomdata_first(cl_atomdata_t *ad, int ntypes, gmx_device_runtime_data_t *runData)
{
cl_int cl_error;
ad->shift_vec_elem_size = sizeof(*(((nbnxn_atomdata_t*)0)->shift_vec));
// TODO: handle errors, check clCreateBuffer flags
- ad->shift_vec = clCreateBuffer(dev_info->context, CL_MEM_READ_WRITE, SHIFTS * ad->shift_vec_elem_size, NULL, &cl_error);
+ ad->shift_vec = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, SHIFTS * ad->shift_vec_elem_size, NULL, &cl_error);
assert(cl_error == CL_SUCCESS);
ad->bShiftVecUploaded = false;
of the host side fshift buffer. */
ad->fshift_elem_size = sizeof(*(((cl_nb_staging_t*)0)->fshift));
- ad->fshift = clCreateBuffer(dev_info->context, CL_MEM_READ_WRITE, SHIFTS * ad->fshift_elem_size, NULL, &cl_error);
+ ad->fshift = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, SHIFTS * ad->fshift_elem_size, NULL, &cl_error);
assert(cl_error == CL_SUCCESS);
// TODO: handle errors, check clCreateBuffer flags
- ad->e_lj = clCreateBuffer(dev_info->context, CL_MEM_READ_WRITE, sizeof(float), NULL, &cl_error);
+ ad->e_lj = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, sizeof(float), NULL, &cl_error);
assert(cl_error == CL_SUCCESS);
// TODO: handle errors, check clCreateBuffer flags
- ad->e_el = clCreateBuffer(dev_info->context, CL_MEM_READ_WRITE, sizeof(float), NULL, &cl_error);
+ ad->e_el = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, sizeof(float), NULL, &cl_error);
assert(cl_error == CL_SUCCESS);
// TODO: handle errors, check clCreateBuffer flags
* evdwOcl. */
static void
map_interaction_types_to_gpu_kernel_flavors(const interaction_const_t *ic,
+ int combRule,
int *gpu_eeltype,
int *gpu_vdwtype)
{
{
case eintmodNONE:
case eintmodPOTSHIFT:
- *gpu_vdwtype = evdwOclCUT;
+ switch (combRule)
+ {
+ case ljcrNONE:
+ *gpu_vdwtype = evdwOclCUT;
+ break;
+ case ljcrGEOM:
+ *gpu_vdwtype = evdwOclCUTCOMBGEOM;
+ break;
+ case ljcrLB:
+ *gpu_vdwtype = evdwOclCUTCOMBLB;
+ break;
+ default:
+ gmx_incons("The requested LJ combination rule is not implemented in the OpenCL GPU accelerated kernels!");
+ break;
+ }
break;
case eintmodFORCESWITCH:
*gpu_vdwtype = evdwOclFSWITCH;
/*! \brief Initializes the nonbonded parameter data structure.
*/
-static void init_nbparam(cl_nbparam_t *nbp,
- const interaction_const_t *ic,
- const nbnxn_atomdata_t *nbat,
- const gmx_device_info_t *dev_info)
+static void init_nbparam(cl_nbparam_t *nbp,
+ const interaction_const_t *ic,
+ const nbnxn_atomdata_t *nbat,
+ const gmx_device_runtime_data_t *runData)
{
int ntypes, nnbfp, nnbfp_comb;
cl_int cl_error;
set_cutoff_parameters(nbp, ic);
map_interaction_types_to_gpu_kernel_flavors(ic,
+ nbat->comb_rule,
&(nbp->eeltype),
&(nbp->vdwtype));
nbp->coulomb_tab_climg2d = NULL;
if (nbp->eeltype == eelOclEWALD_TAB || nbp->eeltype == eelOclEWALD_TAB_TWIN)
{
- init_ewald_coulomb_force_table(ic, nbp, dev_info);
+ init_ewald_coulomb_force_table(ic, nbp, runData);
}
else
// TODO: improvement needed.
array_format.image_channel_data_type = CL_FLOAT;
array_format.image_channel_order = CL_R;
- nbp->coulomb_tab_climg2d = clCreateImage2D(dev_info->context, CL_MEM_READ_WRITE,
+ nbp->coulomb_tab_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE,
&array_format, 1, 1, 0, NULL, &cl_error);
*/
- nbp->coulomb_tab_climg2d = clCreateBuffer(dev_info->context, CL_MEM_READ_ONLY, sizeof(cl_float), NULL, &cl_error);
+ nbp->coulomb_tab_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY, sizeof(cl_float), NULL, &cl_error);
// TODO: handle errors
}
array_format.image_channel_data_type = CL_FLOAT;
array_format.image_channel_order = CL_R;
- nbp->nbfp_climg2d = clCreateImage2D(dev_info->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+ nbp->nbfp_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
&array_format, nnbfp, 1, 0, nbat->nbfp, &cl_error);
*/
- nbp->nbfp_climg2d = clCreateBuffer(dev_info->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, nnbfp*sizeof(cl_float), nbat->nbfp, &cl_error);
+ nbp->nbfp_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, nnbfp*sizeof(cl_float), nbat->nbfp, &cl_error);
assert(cl_error == CL_SUCCESS);
// TODO: handle errors
{
/* Switched from using textures to using buffers */
// TODO: decide which alternative is most efficient - textures or buffers.
- /* nbp->nbfp_comb_climg2d = clCreateImage2D(dev_info->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+ /* nbp->nbfp_comb_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
&array_format, nnbfp_comb, 1, 0, nbat->nbfp_comb, &cl_error);*/
- nbp->nbfp_comb_climg2d = clCreateBuffer(dev_info->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, nnbfp_comb*sizeof(cl_float), nbat->nbfp_comb, &cl_error);
+ nbp->nbfp_comb_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, nnbfp_comb*sizeof(cl_float), nbat->nbfp_comb, &cl_error);
assert(cl_error == CL_SUCCESS);
// don't accept NULL values for image2D parameters.
/* Switched from using textures to using buffers */
// TODO: decide which alternative is most efficient - textures or buffers.
- /* nbp->nbfp_comb_climg2d = clCreateImage2D(dev_info->context, CL_MEM_READ_WRITE,
+ /* nbp->nbfp_comb_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE,
&array_format, 1, 1, 0, NULL, &cl_error);*/
- nbp->nbfp_comb_climg2d = clCreateBuffer(dev_info->context, CL_MEM_READ_ONLY, sizeof(cl_float), NULL, &cl_error);
+ nbp->nbfp_comb_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY, sizeof(cl_float), NULL, &cl_error);
assert(cl_error == CL_SUCCESS);
nbp->eeltype = nbnxn_gpu_pick_ewald_kernel_type(ic->rcoulomb != ic->rvdw);
- init_ewald_coulomb_force_table(ic, nb->nbparam, nb->dev_info);
+ init_ewald_coulomb_force_table(ic, nb->nbparam, nb->dev_rundata);
}
/*! \brief Initializes the pair list data structure.
*
* A fatal error results if creation fails.
*
- * \param[inout] nb Manages OpenCL non-bonded calculations;
- * contexts returned in dev_info members
- * \param[in] rank MPI rank (for error reporting)
+ * \param[inout] runtimeData runtime data including program and context
+ * \param[in] devInfo device info struct
+ * \param[in] rank MPI rank (for error reporting)
*/
static void
-nbnxn_gpu_create_context(gmx_nbnxn_ocl_t *nb,
+nbnxn_gpu_create_context(gmx_device_runtime_data_t *runtimeData,
+ const gmx_device_info_t *devInfo,
int rank)
{
cl_context_properties context_properties[3];
cl_context context;
cl_int cl_error;
- platform_id = nb->dev_info->ocl_gpu_id.ocl_platform_id;
- device_id = nb->dev_info->ocl_gpu_id.ocl_device_id;
+ assert(runtimeData != NULL);
+ assert(devInfo != NULL);
+
+ platform_id = devInfo->ocl_gpu_id.ocl_platform_id;
+ device_id = devInfo->ocl_gpu_id.ocl_device_id;
context_properties[0] = CL_CONTEXT_PLATFORM;
context_properties[1] = (cl_context_properties) platform_id;
context = clCreateContext(context_properties, 1, &device_id, NULL, NULL, &cl_error);
if (CL_SUCCESS != cl_error)
{
- gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s: OpenCL error %d",
+ gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s:\n OpenCL error %d: %s",
rank,
- nb->dev_info->device_name,
- cl_error);
+ devInfo->device_name,
+ cl_error, ocl_get_error_string(cl_error).c_str());
return;
}
- nb->dev_info->context = context;
+ runtimeData->context = context;
}
/*! \brief Initializes the OpenCL kernel pointers of the nbnxn_ocl_ptr_t input data structure. */
cl_kernel kernel;
cl_int cl_error;
- kernel = clCreateKernel(nb->dev_info->program, kernel_name, &cl_error);
+ kernel = clCreateKernel(nb->dev_rundata->program, kernel_name, &cl_error);
if (CL_SUCCESS != cl_error)
{
gmx_fatal(FARGS, "Failed to create kernel '%s' for GPU #%s: OpenCL error %d",
cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_mem), &(adat->e_lj));
cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_mem), &(adat->e_el));
cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_uint), &shifts);
- assert(cl_error == CL_SUCCESS);
+ GMX_ASSERT(cl_error == CL_SUCCESS, ocl_get_error_string(cl_error).c_str());
cl_error = clEnqueueNDRangeKernel(ls, zero_e_fshift, 3, NULL, global_work_size, local_work_size, 0, NULL, NULL);
- assert(cl_error == CL_SUCCESS);
-
+ GMX_ASSERT(cl_error == CL_SUCCESS, ocl_get_error_string(cl_error).c_str());
}
/*! \brief Initializes the OpenCL kernel pointers of the nbnxn_ocl_ptr_t input data structure. */
nb->kernel_zero_e_fshift = nbnxn_gpu_create_kernel(nb, "zero_e_fshift");
}
+/*! \brief Initializes simulation constant data.
+ *
+ * Initializes members of the atomdata and nbparam structs and
+ * clears e/fshift output buffers.
+ */
+static void nbnxn_ocl_init_const(gmx_nbnxn_ocl_t *nb,
+ const interaction_const_t *ic,
+ const nonbonded_verlet_group_t *nbv_group)
+{
+ init_atomdata_first(nb->atdat, nbv_group[0].nbat->ntype, nb->dev_rundata);
+ init_nbparam(nb->nbparam, ic, nbv_group[0].nbat, nb->dev_rundata);
+}
+
+
//! This function is documented in the header file
-void nbnxn_gpu_init(FILE gmx_unused *fplog,
- gmx_nbnxn_ocl_t **p_nb,
+void nbnxn_gpu_init(gmx_nbnxn_ocl_t **p_nb,
const gmx_gpu_info_t *gpu_info,
const gmx_gpu_opt_t *gpu_opt,
const interaction_const_t *ic,
{
gmx_nbnxn_ocl_t *nb;
cl_int cl_error;
- /*
- bool gmx_unused bStreamSync;
- bool gmx_unused bNoStreamSync;
- bool gmx_unused bTMPIAtomics;
- bool gmx_unused bX86;
- bool gmx_unused bOldDriver;
- */
cl_command_queue_properties queue_properties;
assert(gpu_info);
/* set device info, just point it to the right GPU among the detected ones */
nb->dev_info = gpu_info->gpu_dev + gpu_opt->dev_use[my_gpu_index];
+ snew(nb->dev_rundata, 1);
/* init to NULL the debug buffer */
nb->debug_buffer = NULL;
queue_properties = 0;
}
- nbnxn_gpu_create_context(nb, rank);
+ nbnxn_gpu_create_context(nb->dev_rundata, nb->dev_info, rank);
/* local/non-local GPU streams */
- nb->stream[eintLocal] = clCreateCommandQueue(nb->dev_info->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
+ nb->stream[eintLocal] = clCreateCommandQueue(nb->dev_rundata->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
if (CL_SUCCESS != cl_error)
{
gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s: OpenCL error %d",
{
init_plist(nb->plist[eintNonlocal]);
- nb->stream[eintNonlocal] = clCreateCommandQueue(nb->dev_info->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
+ nb->stream[eintNonlocal] = clCreateCommandQueue(nb->dev_rundata->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
if (CL_SUCCESS != cl_error)
{
gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s: OpenCL error %d",
init_timings(nb->timings);
}
- // TODO: check if it's worth implementing for NVIDIA GPUs
- ///////////* set the kernel type for the current GPU */
- ///////////* pick L1 cache configuration */
- //////////nbnxn_gpu_set_cacheconfig(nb->dev_info);
+ nbnxn_ocl_init_const(nb, ic, nbv_grp);
+
+ /* Enable LJ param manual prefetch for AMD or if we request through env. var.
+ * TODO: decide about NVIDIA
+ */
+ nb->bPrefetchLjParam =
+ (getenv("GMX_OCL_DISABLE_I_PREFETCH") == NULL) &&
+ ((nb->dev_info->vendor_e == OCL_VENDOR_AMD) || (getenv("GMX_OCL_ENABLE_I_PREFETCH") != NULL));
- init_atomdata_first(nb->atdat, nbv_grp[0].nbat->ntype, nb->dev_info);
- init_nbparam(nb->nbparam, ic, nbv_grp[0].nbat, nb->dev_info);
+ /* NOTE: in CUDA we pick L1 cache configuration for the nbnxn kernels here,
+ * but sadly this is not supported in OpenCL (yet?). Consider adding it if
+ * it becomes supported.
+ */
nbnxn_gpu_compile_kernels(nb);
nbnxn_gpu_init_kernels(nb);
- // TODO put this elsewhere? also mirror it in cuda
+
+ /* clear energy and shift force outputs */
nbnxn_ocl_clear_e_fshift(nb);
*p_nb = nb;
ocl_realloc_buffered(&d_plist->sci, h_plist->sci, sizeof(nbnxn_sci_t),
&d_plist->nsci, &d_plist->sci_nalloc,
h_plist->nsci,
- nb->dev_info->context,
+ nb->dev_rundata->context,
stream, true, &(nb->timers->pl_h2d_sci[iloc]));
ocl_realloc_buffered(&d_plist->cj4, h_plist->cj4, sizeof(nbnxn_cj4_t),
&d_plist->ncj4, &d_plist->cj4_nalloc,
h_plist->ncj4,
- nb->dev_info->context,
+ nb->dev_rundata->context,
stream, true, &(nb->timers->pl_h2d_cj4[iloc]));
ocl_realloc_buffered(&d_plist->excl, h_plist->excl, sizeof(nbnxn_excl_t),
&d_plist->nexcl, &d_plist->excl_nalloc,
h_plist->nexcl,
- nb->dev_info->context,
+ nb->dev_rundata->context,
stream, true, &(nb->timers->pl_h2d_excl[iloc]));
/* need to prune the pair list during the next step */
{
ocl_free_buffered(d_atdat->f, &d_atdat->natoms, &d_atdat->nalloc);
ocl_free_buffered(d_atdat->xq, NULL, NULL);
+ ocl_free_buffered(d_atdat->lj_comb, NULL, NULL);
ocl_free_buffered(d_atdat->atom_types, NULL, NULL);
}
d_atdat->f_elem_size = sizeof(rvec);
// TODO: handle errors, check clCreateBuffer flags
- d_atdat->f = clCreateBuffer(nb->dev_info->context, CL_MEM_READ_WRITE, nalloc * d_atdat->f_elem_size, NULL, &cl_error);
+ d_atdat->f = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * d_atdat->f_elem_size, NULL, &cl_error);
assert(CL_SUCCESS == cl_error);
- d_atdat->xq = clCreateBuffer(nb->dev_info->context, CL_MEM_READ_WRITE, nalloc * sizeof(cl_float4), NULL, &cl_error);
+ // TODO: change the flag to read-only
+ d_atdat->xq = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(cl_float4), NULL, &cl_error);
assert(CL_SUCCESS == cl_error);
// TODO: handle errors, check clCreateBuffer flags
- d_atdat->atom_types = clCreateBuffer(nb->dev_info->context, CL_MEM_READ_WRITE, nalloc * sizeof(int), NULL, &cl_error);
- assert(CL_SUCCESS == cl_error);
- // TODO: handle errors, check clCreateBuffer flags
+ if (useLjCombRule(nb->nbparam->vdwtype))
+ {
+ // TODO: change the flag to read-only
+ d_atdat->lj_comb = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(cl_float2), NULL, &cl_error);
+ assert(CL_SUCCESS == cl_error);
+ // TODO: handle errors, check clCreateBuffer flags
+ }
+ else
+ {
+ // TODO: change the flag to read-only
+ d_atdat->atom_types = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(int), NULL, &cl_error);
+ assert(CL_SUCCESS == cl_error);
+ // TODO: handle errors, check clCreateBuffer flags
+ }
d_atdat->nalloc = nalloc;
realloced = true;
nbnxn_ocl_clear_f(nb, nalloc);
}
- ocl_copy_H2D_async(d_atdat->atom_types, nbat->type, 0,
- natoms*sizeof(int), ls, bDoTime ? &(timers->atdat) : NULL);
+ if (useLjCombRule(nb->nbparam->vdwtype))
+ {
+ ocl_copy_H2D_async(d_atdat->lj_comb, nbat->lj_comb, 0,
+ natoms*sizeof(cl_float2), ls, bDoTime ? &(timers->atdat) : NULL);
+ }
+ else
+ {
+ ocl_copy_H2D_async(d_atdat->atom_types, nbat->type, 0,
+ natoms*sizeof(int), ls, bDoTime ? &(timers->atdat) : NULL);
+
+ }
/* kick off the tasks enqueued above to ensure concurrency with the search */
cl_error = clFlush(ls);
}
}
+/*! \brief Free the OpenCL runtime data (context and program).
+ *
+ * The function releases the OpenCL context and program assuciated with the
+ * device that the calling PP rank is running on.
+ *
+ * \param runData [in] porinter to the structure with runtime data.
+ */
+static void free_gpu_device_runtime_data(gmx_device_runtime_data_t *runData)
+{
+ if (runData == NULL)
+ {
+ return;
+ }
+
+ cl_int gmx_unused cl_error;
+
+ if (runData->context)
+ {
+ cl_error = clReleaseContext(runData->context);
+ runData->context = NULL;
+ assert(CL_SUCCESS == cl_error);
+ }
+
+ if (runData->program)
+ {
+ cl_error = clReleaseProgram(runData->program);
+ runData->program = NULL;
+ assert(CL_SUCCESS == cl_error);
+ }
+
+}
+
//! This function is documented in the header file
void nbnxn_gpu_free(gmx_nbnxn_ocl_t *nb)
{
free_ocl_buffer(&(nb->atdat->e_lj));
free_ocl_buffer(&(nb->atdat->e_el));
free_ocl_buffer(&(nb->atdat->fshift));
+ free_ocl_buffer(&(nb->atdat->lj_comb));
free_ocl_buffer(&(nb->atdat->atom_types));
free_ocl_buffer(&(nb->atdat->shift_vec));
sfree(nb->atdat);
nb->misc_ops_and_local_H2D_done = NULL;
}
+ free_gpu_device_runtime_data(nb->dev_rundata);
+ sfree(nb->dev_rundata);
+
/* Free timers and timings */
sfree(nb->timers);
sfree(nb->timings);
}
}
+
//! This function is documented in the header file
gmx_wallclock_gpu_t * nbnxn_gpu_get_timings(gmx_nbnxn_ocl_t *nb)
{
biod.pnpi.spb.ru / alexxy / gromacs.git / blobdiff