src/gromacs/mdlib/nbnxn_ocl/nbnxn_ocl_data_mgmt.cpp

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2012,2013,2014,2015,2016,2017,2018, by the GROMACS development team, led by
   5  * Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   6  * and including many others, as listed in the AUTHORS file in the
   7  * top-level source directory and at http://www.gromacs.org.
   8  *
   9  * GROMACS is free software; you can redistribute it and/or
  10  * modify it under the terms of the GNU Lesser General Public License
  11  * as published by the Free Software Foundation; either version 2.1
  12  * of the License, or (at your option) any later version.
  13  *
  14  * GROMACS is distributed in the hope that it will be useful,
  15  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  16  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  17  * Lesser General Public License for more details.
  18  *
  19  * You should have received a copy of the GNU Lesser General Public
  20  * License along with GROMACS; if not, see
  21  * http://www.gnu.org/licenses, or write to the Free Software Foundation,
  22  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  23  *
  24  * If you want to redistribute modifications to GROMACS, please
  25  * consider that scientific software is very special. Version
  26  * control is crucial - bugs must be traceable. We will be happy to
  27  * consider code for inclusion in the official distribution, but
  28  * derived work must not be called official GROMACS. Details are found
  29  * in the README & COPYING files - if they are missing, get the
  30  * official version at http://www.gromacs.org.
  31  *
  32  * To help us fund GROMACS development, we humbly ask that you cite
  33  * the research papers on the package. Check out http://www.gromacs.org.
  34  */
  35 /*! \internal \file
  36  *  \brief Define OpenCL implementation of nbnxn_gpu_data_mgmt.h
  37  *
  38  *  \author Anca Hamuraru <anca@streamcomputing.eu>
  39  *  \author Dimitrios Karkoulis <dimitris.karkoulis@gmail.com>
  40  *  \author Teemu Virolainen <teemu@streamcomputing.eu>
  41  *  \author Szilárd Páll <pall.szilard@gmail.com>
  42  */
  43 #include "gmxpre.h"
  44
  45 #include <assert.h>
  46 #include <stdarg.h>
  47 #include <stdio.h>
  48 #include <stdlib.h>
  49 #include <string.h>
  50
  51 #include <cmath>
  52
  53 #include "gromacs/gpu_utils/gpu_utils.h"
  54 #include "gromacs/gpu_utils/oclutils.h"
  55 #include "gromacs/hardware/gpu_hw_info.h"
  56 #include "gromacs/math/vectypes.h"
  57 #include "gromacs/mdlib/force_flags.h"
  58 #include "gromacs/mdlib/nb_verlet.h"
  59 #include "gromacs/mdlib/nbnxn_consts.h"
  60 #include "gromacs/mdlib/nbnxn_gpu.h"
  61 #include "gromacs/mdlib/nbnxn_gpu_data_mgmt.h"
  62 #include "gromacs/mdlib/nbnxn_gpu_jit_support.h"
  63 #include "gromacs/mdtypes/interaction_const.h"
  64 #include "gromacs/mdtypes/md_enums.h"
  65 #include "gromacs/pbcutil/ishift.h"
  66 #include "gromacs/timing/gpu_timing.h"
  67 #include "gromacs/utility/cstringutil.h"
  68 #include "gromacs/utility/fatalerror.h"
  69 #include "gromacs/utility/gmxassert.h"
  70 #include "gromacs/utility/real.h"
  71 #include "gromacs/utility/smalloc.h"
  72
  73 #include "nbnxn_ocl_internal.h"
  74 #include "nbnxn_ocl_types.h"
  75
  76 /*! \brief This parameter should be determined heuristically from the
  77  * kernel execution times
  78  *
  79  * This value is best for small systems on a single AMD Radeon R9 290X
  80  * (and about 5% faster than 40, which is the default for CUDA
  81  * devices). Larger simulation systems were quite insensitive to the
  82  * value of this parameter.
  83  */
  84 static unsigned int gpu_min_ci_balanced_factor = 50;
  85
  86
  87 /*! \brief Returns true if LJ combination rules are used in the non-bonded kernels.
  88  *
  89  * Full doc in nbnxn_ocl_internal.h */
  90 bool useLjCombRule(int vdwType)
  91 {
  92     return (vdwType == evdwOclCUTCOMBGEOM ||
  93             vdwType == evdwOclCUTCOMBLB);
  94 }
  95
  96 /*! \brief Tabulates the Ewald Coulomb force and initializes the size/scale
  97  * and the table GPU array.
  98  *
  99  * If called with an already allocated table, it just re-uploads the
 100  * table.
 101  */
 102 static void init_ewald_coulomb_force_table(const interaction_const_t       *ic,
 103                                            cl_nbparam_t                    *nbp,
 104                                            const gmx_device_runtime_data_t *runData)
 105 {
 106     cl_mem       coul_tab;
 107
 108     cl_int       cl_error;
 109
 110     if (nbp->coulomb_tab_climg2d != nullptr)
 111     {
 112         freeDeviceBuffer(&(nbp->coulomb_tab_climg2d));
 113     }
 114
 115     /* Switched from using textures to using buffers */
 116     // TODO: decide which alternative is most efficient - textures or buffers.
 117     /*
 118        cl_image_format array_format;
 119
 120        array_format.image_channel_data_type = CL_FLOAT;
 121        array_format.image_channel_order     = CL_R;
 122
 123        coul_tab = clCreateImage2D(runData->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
 124        &array_format, tabsize, 1, 0, ftmp, &cl_error);
 125      */
 126
 127     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);
 128     assert(cl_error == CL_SUCCESS);
 129     // TODO: handle errors, check clCreateBuffer flags
 130
 131     nbp->coulomb_tab_climg2d  = coul_tab;
 132     nbp->coulomb_tab_scale    = ic->tabq_scale;
 133 }
 134
 135
 136 /*! \brief Initializes the atomdata structure first time, it only gets filled at
 137     pair-search.
 138  */
 139 static void init_atomdata_first(cl_atomdata_t *ad, int ntypes, gmx_device_runtime_data_t *runData)
 140 {
 141     cl_int cl_error;
 142
 143     ad->ntypes  = ntypes;
 144
 145     /* An element of the shift_vec device buffer has the same size as one element
 146        of the host side shift_vec buffer. */
 147     ad->shift_vec_elem_size = sizeof(*nbnxn_atomdata_t::shift_vec);
 148
 149     // TODO: handle errors, check clCreateBuffer flags
 150     ad->shift_vec = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, SHIFTS * ad->shift_vec_elem_size, nullptr, &cl_error);
 151     assert(cl_error == CL_SUCCESS);
 152     ad->bShiftVecUploaded = false;
 153
 154     /* An element of the fshift device buffer has the same size as one element
 155        of the host side fshift buffer. */
 156     ad->fshift_elem_size = sizeof(*cl_nb_staging_t::fshift);
 157
 158     ad->fshift = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, SHIFTS * ad->fshift_elem_size, nullptr, &cl_error);
 159     assert(cl_error == CL_SUCCESS);
 160     // TODO: handle errors, check clCreateBuffer flags
 161
 162     ad->e_lj = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, sizeof(float), nullptr, &cl_error);
 163     assert(cl_error == CL_SUCCESS);
 164     // TODO: handle errors, check clCreateBuffer flags
 165
 166     ad->e_el = clCreateBuffer(runData->context, CL_MEM_READ_WRITE, sizeof(float), nullptr, &cl_error);
 167     assert(cl_error == CL_SUCCESS);
 168     // TODO: handle errors, check clCreateBuffer flags
 169
 170     /* initialize to nullptr pointers to data that is not allocated here and will
 171        need reallocation in nbnxn_gpu_init_atomdata */
 172     ad->xq = nullptr;
 173     ad->f  = nullptr;
 174
 175     /* size -1 indicates that the respective array hasn't been initialized yet */
 176     ad->natoms = -1;
 177     ad->nalloc = -1;
 178 }
 179
 180 /*! \brief Copies all parameters related to the cut-off from ic to nbp
 181  */
 182 static void set_cutoff_parameters(cl_nbparam_t              *nbp,
 183                                   const interaction_const_t *ic,
 184                                   const NbnxnListParameters *listParams)
 185 {
 186     nbp->ewald_beta        = ic->ewaldcoeff_q;
 187     nbp->sh_ewald          = ic->sh_ewald;
 188     nbp->epsfac            = ic->epsfac;
 189     nbp->two_k_rf          = 2.0 * ic->k_rf;
 190     nbp->c_rf              = ic->c_rf;
 191     nbp->rvdw_sq           = ic->rvdw * ic->rvdw;
 192     nbp->rcoulomb_sq       = ic->rcoulomb * ic->rcoulomb;
 193     nbp->rlistOuter_sq     = listParams->rlistOuter * listParams->rlistOuter;
 194     nbp->rlistInner_sq     = listParams->rlistInner * listParams->rlistInner;
 195     nbp->useDynamicPruning = listParams->useDynamicPruning;
 196
 197     nbp->sh_lj_ewald       = ic->sh_lj_ewald;
 198     nbp->ewaldcoeff_lj     = ic->ewaldcoeff_lj;
 199
 200     nbp->rvdw_switch       = ic->rvdw_switch;
 201     nbp->dispersion_shift  = ic->dispersion_shift;
 202     nbp->repulsion_shift   = ic->repulsion_shift;
 203     nbp->vdw_switch        = ic->vdw_switch;
 204 }
 205
 206 /*! \brief Returns the kinds of electrostatics and Vdw OpenCL
 207  *  kernels that will be used.
 208  *
 209  * Respectively, these values are from enum eelOcl and enum
 210  * evdwOcl. */
 211 static void
 212 map_interaction_types_to_gpu_kernel_flavors(const interaction_const_t *ic,
 213                                             int                        combRule,
 214                                             int                       *gpu_eeltype,
 215                                             int                       *gpu_vdwtype)
 216 {
 217     if (ic->vdwtype == evdwCUT)
 218     {
 219         switch (ic->vdw_modifier)
 220         {
 221             case eintmodNONE:
 222             case eintmodPOTSHIFT:
 223                 switch (combRule)
 224                 {
 225                     case ljcrNONE:
 226                         *gpu_vdwtype = evdwOclCUT;
 227                         break;
 228                     case ljcrGEOM:
 229                         *gpu_vdwtype = evdwOclCUTCOMBGEOM;
 230                         break;
 231                     case ljcrLB:
 232                         *gpu_vdwtype = evdwOclCUTCOMBLB;
 233                         break;
 234                     default:
 235                         gmx_incons("The requested LJ combination rule is not implemented in the OpenCL GPU accelerated kernels!");
 236                         break;
 237                 }
 238                 break;
 239             case eintmodFORCESWITCH:
 240                 *gpu_vdwtype = evdwOclFSWITCH;
 241                 break;
 242             case eintmodPOTSWITCH:
 243                 *gpu_vdwtype = evdwOclPSWITCH;
 244                 break;
 245             default:
 246                 gmx_incons("The requested VdW interaction modifier is not implemented in the GPU accelerated kernels!");
 247                 break;
 248         }
 249     }
 250     else if (ic->vdwtype == evdwPME)
 251     {
 252         if (ic->ljpme_comb_rule == ljcrGEOM)
 253         {
 254             *gpu_vdwtype = evdwOclEWALDGEOM;
 255         }
 256         else
 257         {
 258             *gpu_vdwtype = evdwOclEWALDLB;
 259         }
 260     }
 261     else
 262     {
 263         gmx_incons("The requested VdW type is not implemented in the GPU accelerated kernels!");
 264     }
 265
 266     if (ic->eeltype == eelCUT)
 267     {
 268         *gpu_eeltype = eelOclCUT;
 269     }
 270     else if (EEL_RF(ic->eeltype))
 271     {
 272         *gpu_eeltype = eelOclRF;
 273     }
 274     else if ((EEL_PME(ic->eeltype) || ic->eeltype == eelEWALD))
 275     {
 276         /* Initially rcoulomb == rvdw, so it's surely not twin cut-off. */
 277         *gpu_eeltype = nbnxn_gpu_pick_ewald_kernel_type(false);
 278     }
 279     else
 280     {
 281         /* Shouldn't happen, as this is checked when choosing Verlet-scheme */
 282         gmx_incons("The requested electrostatics type is not implemented in the GPU accelerated kernels!");
 283     }
 284 }
 285
 286 /*! \brief Initializes the nonbonded parameter data structure.
 287  */
 288 static void init_nbparam(cl_nbparam_t                    *nbp,
 289                          const interaction_const_t       *ic,
 290                          const NbnxnListParameters       *listParams,
 291                          const nbnxn_atomdata_t          *nbat,
 292                          const gmx_device_runtime_data_t *runData)
 293 {
 294     int         ntypes, nnbfp, nnbfp_comb;
 295     cl_int      cl_error;
 296
 297
 298     ntypes = nbat->ntype;
 299
 300     set_cutoff_parameters(nbp, ic, listParams);
 301
 302     map_interaction_types_to_gpu_kernel_flavors(ic,
 303                                                 nbat->comb_rule,
 304                                                 &(nbp->eeltype),
 305                                                 &(nbp->vdwtype));
 306
 307     if (ic->vdwtype == evdwPME)
 308     {
 309         if (ic->ljpme_comb_rule == ljcrGEOM)
 310         {
 311             assert(nbat->comb_rule == ljcrGEOM);
 312         }
 313         else
 314         {
 315             assert(nbat->comb_rule == ljcrLB);
 316         }
 317     }
 318     /* generate table for PME */
 319     nbp->coulomb_tab_climg2d = nullptr;
 320     if (nbp->eeltype == eelOclEWALD_TAB || nbp->eeltype == eelOclEWALD_TAB_TWIN)
 321     {
 322         init_ewald_coulomb_force_table(ic, nbp, runData);
 323     }
 324     else
 325     // TODO: improvement needed.
 326     // The image2d is created here even if eeltype is not eelCuEWALD_TAB or eelCuEWALD_TAB_TWIN because the OpenCL kernels
 327     // don't accept nullptr values for image2D parameters.
 328     {
 329         /* Switched from using textures to using buffers */
 330         // TODO: decide which alternative is most efficient - textures or buffers.
 331         /*
 332            cl_image_format array_format;
 333
 334            array_format.image_channel_data_type = CL_FLOAT;
 335            array_format.image_channel_order     = CL_R;
 336
 337            nbp->coulomb_tab_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE,
 338             &array_format, 1, 1, 0, nullptr, &cl_error);
 339          */
 340
 341         nbp->coulomb_tab_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY, sizeof(cl_float), nullptr, &cl_error);
 342         // TODO: handle errors
 343     }
 344
 345     nnbfp      = 2*ntypes*ntypes;
 346     nnbfp_comb = 2*ntypes;
 347
 348     {
 349         /* Switched from using textures to using buffers */
 350         // TODO: decide which alternative is most efficient - textures or buffers.
 351         /*
 352            cl_image_format array_format;
 353
 354            array_format.image_channel_data_type = CL_FLOAT;
 355            array_format.image_channel_order     = CL_R;
 356
 357            nbp->nbfp_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
 358             &array_format, nnbfp, 1, 0, nbat->nbfp, &cl_error);
 359          */
 360
 361         nbp->nbfp_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, nnbfp*sizeof(cl_float), nbat->nbfp, &cl_error);
 362         assert(cl_error == CL_SUCCESS);
 363         // TODO: handle errors
 364
 365         if (ic->vdwtype == evdwPME)
 366         {
 367             /* Switched from using textures to using buffers */
 368             // TODO: decide which alternative is most efficient - textures or buffers.
 369             /*  nbp->nbfp_comb_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
 370                 &array_format, nnbfp_comb, 1, 0, nbat->nbfp_comb, &cl_error);*/
 371             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);
 372
 373
 374             assert(cl_error == CL_SUCCESS);
 375             // TODO: handle errors
 376         }
 377         else
 378         {
 379             // TODO: improvement needed.
 380             // The image2d is created here even if vdwtype is not evdwPME because the OpenCL kernels
 381             // don't accept nullptr values for image2D parameters.
 382             /* Switched from using textures to using buffers */
 383             // TODO: decide which alternative is most efficient - textures or buffers.
 384             /* nbp->nbfp_comb_climg2d = clCreateImage2D(runData->context, CL_MEM_READ_WRITE,
 385                 &array_format, 1, 1, 0, nullptr, &cl_error);*/
 386             nbp->nbfp_comb_climg2d = clCreateBuffer(runData->context, CL_MEM_READ_ONLY, sizeof(cl_float), nullptr, &cl_error);
 387
 388
 389             assert(cl_error == CL_SUCCESS);
 390             // TODO: handle errors
 391         }
 392     }
 393 }
 394
 395 //! This function is documented in the header file
 396 void nbnxn_gpu_pme_loadbal_update_param(const nonbonded_verlet_t    *nbv,
 397                                         const interaction_const_t   *ic,
 398                                         const NbnxnListParameters   *listParams)
 399 {
 400     if (!nbv || nbv->grp[0].kernel_type != nbnxnk8x8x8_GPU)
 401     {
 402         return;
 403     }
 404     gmx_nbnxn_ocl_t    *nb  = nbv->gpu_nbv;
 405     cl_nbparam_t       *nbp = nb->nbparam;
 406
 407     set_cutoff_parameters(nbp, ic, listParams);
 408
 409     nbp->eeltype = nbnxn_gpu_pick_ewald_kernel_type(ic->rcoulomb != ic->rvdw);
 410
 411     init_ewald_coulomb_force_table(ic, nb->nbparam, nb->dev_rundata);
 412 }
 413
 414 /*! \brief Initializes the pair list data structure.
 415  */
 416 static void init_plist(cl_plist_t *pl)
 417 {
 418     /* initialize to nullptr pointers to data that is not allocated here and will
 419        need reallocation in nbnxn_gpu_init_pairlist */
 420     pl->sci     = nullptr;
 421     pl->cj4     = nullptr;
 422     pl->imask   = nullptr;
 423     pl->excl    = nullptr;
 424
 425     /* size -1 indicates that the respective array hasn't been initialized yet */
 426     pl->na_c           = -1;
 427     pl->nsci           = -1;
 428     pl->sci_nalloc     = -1;
 429     pl->ncj4           = -1;
 430     pl->cj4_nalloc     = -1;
 431     pl->nimask         = -1;
 432     pl->imask_nalloc   = -1;
 433     pl->nexcl          = -1;
 434     pl->excl_nalloc    = -1;
 435     pl->haveFreshList  = false;
 436 }
 437
 438 /*! \brief Initializes the timer data structure.
 439  */
 440 static void init_timers(cl_timers_t *t,
 441                         bool         bUseTwoStreams)
 442 {
 443     for (int i = 0; i <= (bUseTwoStreams ? 1 : 0); i++)
 444     {
 445         t->didPairlistH2D[i]  = false;
 446         t->didPrune[i]        = false;
 447         t->didRollingPrune[i] = false;
 448     }
 449 }
 450
 451 /*! \brief Initializes the timings data structure.
 452  */
 453 static void init_timings(gmx_wallclock_gpu_nbnxn_t *t)
 454 {
 455     int i, j;
 456
 457     t->nb_h2d_t = 0.0;
 458     t->nb_d2h_t = 0.0;
 459     t->nb_c     = 0;
 460     t->pl_h2d_t = 0.0;
 461     t->pl_h2d_c = 0;
 462     for (i = 0; i < 2; i++)
 463     {
 464         for (j = 0; j < 2; j++)
 465         {
 466             t->ktime[i][j].t = 0.0;
 467             t->ktime[i][j].c = 0;
 468         }
 469     }
 470
 471     t->pruneTime.c        = 0;
 472     t->pruneTime.t        = 0.0;
 473     t->dynamicPruneTime.c = 0;
 474     t->dynamicPruneTime.t = 0.0;
 475 }
 476
 477 /*! \brief Creates context for OpenCL GPU given by \p mygpu
 478  *
 479  * A fatal error results if creation fails.
 480  *
 481  * \param[inout] runtimeData runtime data including program and context
 482  * \param[in]    devInfo     device info struct
 483  * \param[in]    rank        MPI rank (for error reporting)
 484  */
 485 static void
 486 nbnxn_gpu_create_context(gmx_device_runtime_data_t *runtimeData,
 487                          const gmx_device_info_t   *devInfo,
 488                          int                        rank)
 489 {
 490     cl_context_properties     context_properties[3];
 491     cl_platform_id            platform_id;
 492     cl_device_id              device_id;
 493     cl_context                context;
 494     cl_int                    cl_error;
 495
 496     assert(runtimeData != nullptr);
 497     assert(devInfo != nullptr);
 498
 499     platform_id      = devInfo->ocl_gpu_id.ocl_platform_id;
 500     device_id        = devInfo->ocl_gpu_id.ocl_device_id;
 501
 502     context_properties[0] = CL_CONTEXT_PLATFORM;
 503     context_properties[1] = (cl_context_properties) platform_id;
 504     context_properties[2] = 0; /* Terminates the list of properties */
 505
 506     context = clCreateContext(context_properties, 1, &device_id, nullptr, nullptr, &cl_error);
 507     if (CL_SUCCESS != cl_error)
 508     {
 509         gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s:\n OpenCL error %d: %s",
 510                   rank,
 511                   devInfo->device_name,
 512                   cl_error, ocl_get_error_string(cl_error).c_str());
 513         return;
 514     }
 515
 516     runtimeData->context = context;
 517 }
 518
 519 /*! \brief Initializes the OpenCL kernel pointers of the nbnxn_ocl_ptr_t input data structure. */
 520 static cl_kernel nbnxn_gpu_create_kernel(gmx_nbnxn_ocl_t *nb,
 521                                          const char      *kernel_name)
 522 {
 523     cl_kernel kernel;
 524     cl_int    cl_error;
 525
 526     kernel = clCreateKernel(nb->dev_rundata->program, kernel_name, &cl_error);
 527     if (CL_SUCCESS != cl_error)
 528     {
 529         gmx_fatal(FARGS, "Failed to create kernel '%s' for GPU #%s: OpenCL error %d",
 530                   kernel_name,
 531                   nb->dev_info->device_name,
 532                   cl_error);
 533     }
 534
 535     return kernel;
 536 }
 537
 538 /*! \brief Clears nonbonded shift force output array and energy outputs on the GPU.
 539  */
 540 static void
 541 nbnxn_ocl_clear_e_fshift(gmx_nbnxn_ocl_t *nb)
 542 {
 543
 544     cl_int               cl_error;
 545     cl_atomdata_t *      adat     = nb->atdat;
 546     cl_command_queue     ls       = nb->stream[eintLocal];
 547
 548     size_t               local_work_size[3]   = {1, 1, 1};
 549     size_t               global_work_size[3]  = {1, 1, 1};
 550
 551     cl_int               shifts   = SHIFTS*3;
 552
 553     cl_int               arg_no;
 554
 555     cl_kernel            zero_e_fshift = nb->kernel_zero_e_fshift;
 556
 557     local_work_size[0]   = 64;
 558     // Round the total number of threads up from the array size
 559     global_work_size[0]  = ((shifts + local_work_size[0] - 1)/local_work_size[0])*local_work_size[0];
 560
 561     arg_no    = 0;
 562     cl_error  = clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_mem), &(adat->fshift));
 563     cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_mem), &(adat->e_lj));
 564     cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_mem), &(adat->e_el));
 565     cl_error |= clSetKernelArg(zero_e_fshift, arg_no++, sizeof(cl_uint), &shifts);
 566     GMX_ASSERT(cl_error == CL_SUCCESS, ocl_get_error_string(cl_error).c_str());
 567
 568     cl_error = clEnqueueNDRangeKernel(ls, zero_e_fshift, 3, nullptr, global_work_size, local_work_size, 0, nullptr, nullptr);
 569     GMX_ASSERT(cl_error == CL_SUCCESS, ocl_get_error_string(cl_error).c_str());
 570 }
 571
 572 /*! \brief Initializes the OpenCL kernel pointers of the nbnxn_ocl_ptr_t input data structure. */
 573 static void nbnxn_gpu_init_kernels(gmx_nbnxn_ocl_t *nb)
 574 {
 575     /* Init to 0 main kernel arrays */
 576     /* They will be later on initialized in select_nbnxn_kernel */
 577     // TODO: consider always creating all variants of the kernels here so that there is no
 578     // need for late call to clCreateKernel -- if that gives any advantage?
 579     memset(nb->kernel_ener_noprune_ptr, 0, sizeof(nb->kernel_ener_noprune_ptr));
 580     memset(nb->kernel_ener_prune_ptr, 0, sizeof(nb->kernel_ener_prune_ptr));
 581     memset(nb->kernel_noener_noprune_ptr, 0, sizeof(nb->kernel_noener_noprune_ptr));
 582     memset(nb->kernel_noener_prune_ptr, 0, sizeof(nb->kernel_noener_prune_ptr));
 583
 584     /* Init pruning kernels
 585      *
 586      * TODO: we could avoid creating kernels if dynamic pruning is turned off,
 587      * but ATM that depends on force flags not passed into the initialization.
 588      */
 589     nb->kernel_pruneonly[epruneFirst]   = nbnxn_gpu_create_kernel(nb, "nbnxn_kernel_prune_opencl");
 590     nb->kernel_pruneonly[epruneRolling] = nbnxn_gpu_create_kernel(nb, "nbnxn_kernel_prune_rolling_opencl");
 591
 592     /* Init auxiliary kernels */
 593     nb->kernel_memset_f      = nbnxn_gpu_create_kernel(nb, "memset_f");
 594     nb->kernel_memset_f2     = nbnxn_gpu_create_kernel(nb, "memset_f2");
 595     nb->kernel_memset_f3     = nbnxn_gpu_create_kernel(nb, "memset_f3");
 596     nb->kernel_zero_e_fshift = nbnxn_gpu_create_kernel(nb, "zero_e_fshift");
 597 }
 598
 599 /*! \brief Initializes simulation constant data.
 600  *
 601  *  Initializes members of the atomdata and nbparam structs and
 602  *  clears e/fshift output buffers.
 603  */
 604 static void nbnxn_ocl_init_const(gmx_nbnxn_ocl_t                *nb,
 605                                  const interaction_const_t      *ic,
 606                                  const NbnxnListParameters      *listParams,
 607                                  const nbnxn_atomdata_t         *nbat)
 608 {
 609     init_atomdata_first(nb->atdat, nbat->ntype, nb->dev_rundata);
 610     init_nbparam(nb->nbparam, ic, listParams, nbat, nb->dev_rundata);
 611 }
 612
 613
 614 //! This function is documented in the header file
 615 void nbnxn_gpu_init(gmx_nbnxn_ocl_t          **p_nb,
 616                     const gmx_device_info_t   *deviceInfo,
 617                     const interaction_const_t *ic,
 618                     const NbnxnListParameters *listParams,
 619                     const nbnxn_atomdata_t    *nbat,
 620                     int                        rank,
 621                     gmx_bool                   bLocalAndNonlocal)
 622 {
 623     gmx_nbnxn_ocl_t            *nb;
 624     cl_int                      cl_error;
 625     cl_command_queue_properties queue_properties;
 626
 627     assert(ic);
 628
 629     if (p_nb == nullptr)
 630     {
 631         return;
 632     }
 633
 634     snew(nb, 1);
 635     snew(nb->atdat, 1);
 636     snew(nb->nbparam, 1);
 637     snew(nb->plist[eintLocal], 1);
 638     if (bLocalAndNonlocal)
 639     {
 640         snew(nb->plist[eintNonlocal], 1);
 641     }
 642
 643     nb->bUseTwoStreams = bLocalAndNonlocal;
 644
 645     nb->timers = new cl_timers_t();
 646     snew(nb->timings, 1);
 647
 648     /* set device info, just point it to the right GPU among the detected ones */
 649     nb->dev_info = deviceInfo;
 650     snew(nb->dev_rundata, 1);
 651
 652     /* init nbst */
 653     pmalloc((void**)&nb->nbst.e_lj, sizeof(*nb->nbst.e_lj));
 654     pmalloc((void**)&nb->nbst.e_el, sizeof(*nb->nbst.e_el));
 655     pmalloc((void**)&nb->nbst.fshift, SHIFTS * sizeof(*nb->nbst.fshift));
 656
 657     init_plist(nb->plist[eintLocal]);
 658
 659     /* OpenCL timing disabled if GMX_DISABLE_GPU_TIMING is defined. */
 660     nb->bDoTime = (getenv("GMX_DISABLE_GPU_TIMING") == nullptr);
 661
 662     /* Create queues only after bDoTime has been initialized */
 663     if (nb->bDoTime)
 664     {
 665         queue_properties = CL_QUEUE_PROFILING_ENABLE;
 666     }
 667     else
 668     {
 669         queue_properties = 0;
 670     }
 671
 672     nbnxn_gpu_create_context(nb->dev_rundata, nb->dev_info, rank);
 673
 674     /* local/non-local GPU streams */
 675     nb->stream[eintLocal] = clCreateCommandQueue(nb->dev_rundata->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
 676     if (CL_SUCCESS != cl_error)
 677     {
 678         gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s: OpenCL error %d",
 679                   rank,
 680                   nb->dev_info->device_name,
 681                   cl_error);
 682         return;
 683     }
 684
 685     if (nb->bUseTwoStreams)
 686     {
 687         init_plist(nb->plist[eintNonlocal]);
 688
 689         nb->stream[eintNonlocal] = clCreateCommandQueue(nb->dev_rundata->context, nb->dev_info->ocl_gpu_id.ocl_device_id, queue_properties, &cl_error);
 690         if (CL_SUCCESS != cl_error)
 691         {
 692             gmx_fatal(FARGS, "On rank %d failed to create context for GPU #%s: OpenCL error %d",
 693                       rank,
 694                       nb->dev_info->device_name,
 695                       cl_error);
 696             return;
 697         }
 698     }
 699
 700     if (nb->bDoTime)
 701     {
 702         init_timers(nb->timers, nb->bUseTwoStreams);
 703         init_timings(nb->timings);
 704     }
 705
 706     nbnxn_ocl_init_const(nb, ic, listParams, nbat);
 707
 708     /* Enable LJ param manual prefetch for AMD or if we request through env. var.
 709      * TODO: decide about NVIDIA
 710      */
 711     nb->bPrefetchLjParam =
 712         (getenv("GMX_OCL_DISABLE_I_PREFETCH") == nullptr) &&
 713         ((nb->dev_info->vendor_e == OCL_VENDOR_AMD) || (getenv("GMX_OCL_ENABLE_I_PREFETCH") != nullptr));
 714
 715     /* NOTE: in CUDA we pick L1 cache configuration for the nbnxn kernels here,
 716      * but sadly this is not supported in OpenCL (yet?). Consider adding it if
 717      * it becomes supported.
 718      */
 719     nbnxn_gpu_compile_kernels(nb);
 720     nbnxn_gpu_init_kernels(nb);
 721
 722     /* clear energy and shift force outputs */
 723     nbnxn_ocl_clear_e_fshift(nb);
 724
 725     *p_nb = nb;
 726
 727     if (debug)
 728     {
 729         fprintf(debug, "Initialized OpenCL data structures.\n");
 730     }
 731 }
 732
 733 /*! \brief Clears the first natoms_clear elements of the GPU nonbonded force output array.
 734  */
 735 static void nbnxn_ocl_clear_f(gmx_nbnxn_ocl_t *nb, int natoms_clear)
 736 {
 737     if (natoms_clear == 0)
 738     {
 739         return;
 740     }
 741
 742     cl_atomdata_t *      adat     = nb->atdat;
 743     cl_command_queue     ls       = nb->stream[eintLocal];
 744     cl_float             value    = 0.0f;
 745
 746     size_t               local_work_size[3]  = {1, 1, 1};
 747     size_t               global_work_size[3] = {1, 1, 1};
 748
 749     cl_int               arg_no;
 750
 751     cl_kernel            memset_f = nb->kernel_memset_f;
 752
 753     cl_uint              natoms_flat = natoms_clear * (sizeof(rvec)/sizeof(real));
 754
 755     local_work_size[0]  = 64;
 756     // Round the total number of threads up from the array size
 757     global_work_size[0] = ((natoms_flat + local_work_size[0] - 1)/local_work_size[0])*local_work_size[0];
 758
 759
 760     cl_int gmx_used_in_debug cl_error;
 761     arg_no    = 0;
 762     cl_error  = clSetKernelArg(memset_f, arg_no++, sizeof(cl_mem), &(adat->f));
 763     cl_error |= clSetKernelArg(memset_f, arg_no++, sizeof(cl_float), &value);
 764     cl_error |= clSetKernelArg(memset_f, arg_no++, sizeof(cl_uint), &natoms_flat);
 765     assert(cl_error == CL_SUCCESS);
 766
 767     cl_error = clEnqueueNDRangeKernel(ls, memset_f, 3, nullptr, global_work_size, local_work_size, 0, nullptr, nullptr);
 768     assert(cl_error == CL_SUCCESS);
 769 }
 770
 771 //! This function is documented in the header file
 772 void
 773 nbnxn_gpu_clear_outputs(gmx_nbnxn_ocl_t   *nb,
 774                         int                flags)
 775 {
 776     nbnxn_ocl_clear_f(nb, nb->atdat->natoms);
 777     /* clear shift force array and energies if the outputs were
 778        used in the current step */
 779     if (flags & GMX_FORCE_VIRIAL)
 780     {
 781         nbnxn_ocl_clear_e_fshift(nb);
 782     }
 783
 784     /* kick off buffer clearing kernel to ensure concurrency with constraints/update */
 785     cl_int gmx_unused cl_error;
 786     cl_error = clFlush(nb->stream[eintLocal]);
 787     assert(CL_SUCCESS == cl_error);
 788 }
 789
 790 //! This function is documented in the header file
 791 void nbnxn_gpu_init_pairlist(gmx_nbnxn_ocl_t        *nb,
 792                              const nbnxn_pairlist_t *h_plist,
 793                              int                     iloc)
 794 {
 795     char             sbuf[STRLEN];
 796     // Timing accumulation should happen only if there was work to do
 797     // because getLastRangeTime() gets skipped with empty lists later
 798     // which leads to the counter not being reset.
 799     bool             bDoTime    = (nb->bDoTime && h_plist->nsci > 0);
 800     cl_command_queue stream     = nb->stream[iloc];
 801     cl_plist_t      *d_plist    = nb->plist[iloc];
 802
 803     if (d_plist->na_c < 0)
 804     {
 805         d_plist->na_c = h_plist->na_ci;
 806     }
 807     else
 808     {
 809         if (d_plist->na_c != h_plist->na_ci)
 810         {
 811             sprintf(sbuf, "In cu_init_plist: the #atoms per cell has changed (from %d to %d)",
 812                     d_plist->na_c, h_plist->na_ci);
 813             gmx_incons(sbuf);
 814         }
 815     }
 816
 817     if (bDoTime)
 818     {
 819         nb->timers->pl_h2d[iloc].openTimingRegion(stream);
 820         nb->timers->didPairlistH2D[iloc] = true;
 821     }
 822
 823     // TODO most of this function is same in CUDA and OpenCL, move into the header
 824     Context context = nb->dev_rundata->context;
 825
 826     reallocateDeviceBuffer(&d_plist->sci, h_plist->nsci,
 827                            &d_plist->nsci, &d_plist->sci_nalloc, context);
 828     copyToDeviceBuffer(&d_plist->sci, h_plist->sci, 0, h_plist->nsci,
 829                        stream, GpuApiCallBehavior::Async,
 830                        bDoTime ? nb->timers->pl_h2d[iloc].fetchNextEvent() : nullptr);
 831
 832     reallocateDeviceBuffer(&d_plist->cj4, h_plist->ncj4,
 833                            &d_plist->ncj4, &d_plist->cj4_nalloc, context);
 834     copyToDeviceBuffer(&d_plist->cj4, h_plist->cj4, 0, h_plist->ncj4,
 835                        stream, GpuApiCallBehavior::Async,
 836                        bDoTime ? nb->timers->pl_h2d[iloc].fetchNextEvent() : nullptr);
 837
 838     reallocateDeviceBuffer(&d_plist->imask, h_plist->ncj4*c_nbnxnGpuClusterpairSplit,
 839                            &d_plist->nimask, &d_plist->imask_nalloc, context);
 840
 841     reallocateDeviceBuffer(&d_plist->excl, h_plist->nexcl,
 842                            &d_plist->nexcl, &d_plist->excl_nalloc, context);
 843     copyToDeviceBuffer(&d_plist->excl, h_plist->excl, 0, h_plist->nexcl,
 844                        stream, GpuApiCallBehavior::Async,
 845                        bDoTime ? nb->timers->pl_h2d[iloc].fetchNextEvent() : nullptr);
 846
 847     if (bDoTime)
 848     {
 849         nb->timers->pl_h2d[iloc].closeTimingRegion(stream);
 850     }
 851
 852     /* need to prune the pair list during the next step */
 853     d_plist->haveFreshList = true;
 854 }
 855
 856 //! This function is documented in the header file
 857 void nbnxn_gpu_upload_shiftvec(gmx_nbnxn_ocl_t        *nb,
 858                                const nbnxn_atomdata_t *nbatom)
 859 {
 860     cl_atomdata_t   *adat  = nb->atdat;
 861     cl_command_queue ls    = nb->stream[eintLocal];
 862
 863     /* only if we have a dynamic box */
 864     if (nbatom->bDynamicBox || !adat->bShiftVecUploaded)
 865     {
 866         ocl_copy_H2D_async(adat->shift_vec, nbatom->shift_vec, 0,
 867                            SHIFTS * adat->shift_vec_elem_size, ls, nullptr);
 868         adat->bShiftVecUploaded = true;
 869     }
 870 }
 871
 872 //! This function is documented in the header file
 873 void nbnxn_gpu_init_atomdata(gmx_nbnxn_ocl_t               *nb,
 874                              const nbnxn_atomdata_t        *nbat)
 875 {
 876     cl_int           cl_error;
 877     int              nalloc, natoms;
 878     bool             realloced;
 879     bool             bDoTime = nb->bDoTime;
 880     cl_timers_t     *timers  = nb->timers;
 881     cl_atomdata_t   *d_atdat = nb->atdat;
 882     cl_command_queue ls      = nb->stream[eintLocal];
 883
 884     natoms    = nbat->natoms;
 885     realloced = false;
 886
 887     if (bDoTime)
 888     {
 889         /* time async copy */
 890         timers->atdat.openTimingRegion(ls);
 891     }
 892
 893     /* need to reallocate if we have to copy more atoms than the amount of space
 894        available and only allocate if we haven't initialized yet, i.e d_atdat->natoms == -1 */
 895     if (natoms > d_atdat->nalloc)
 896     {
 897         nalloc = over_alloc_small(natoms);
 898
 899         /* free up first if the arrays have already been initialized */
 900         if (d_atdat->nalloc != -1)
 901         {
 902             freeDeviceBuffer(&d_atdat->f);
 903             freeDeviceBuffer(&d_atdat->xq);
 904             freeDeviceBuffer(&d_atdat->lj_comb);
 905             freeDeviceBuffer(&d_atdat->atom_types);
 906         }
 907
 908         d_atdat->f_elem_size = sizeof(rvec);
 909
 910         // TODO: handle errors, check clCreateBuffer flags
 911         d_atdat->f = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * d_atdat->f_elem_size, nullptr, &cl_error);
 912         assert(CL_SUCCESS == cl_error);
 913
 914         // TODO: change the flag to read-only
 915         d_atdat->xq = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(cl_float4), nullptr, &cl_error);
 916         assert(CL_SUCCESS == cl_error);
 917         // TODO: handle errors, check clCreateBuffer flags
 918
 919         if (useLjCombRule(nb->nbparam->vdwtype))
 920         {
 921             // TODO: change the flag to read-only
 922             d_atdat->lj_comb = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(cl_float2), nullptr, &cl_error);
 923             assert(CL_SUCCESS == cl_error);
 924             // TODO: handle errors, check clCreateBuffer flags
 925         }
 926         else
 927         {
 928             // TODO: change the flag to read-only
 929             d_atdat->atom_types = clCreateBuffer(nb->dev_rundata->context, CL_MEM_READ_WRITE, nalloc * sizeof(int), nullptr, &cl_error);
 930             assert(CL_SUCCESS == cl_error);
 931             // TODO: handle errors, check clCreateBuffer flags
 932         }
 933
 934         d_atdat->nalloc = nalloc;
 935         realloced       = true;
 936     }
 937
 938     d_atdat->natoms       = natoms;
 939     d_atdat->natoms_local = nbat->natoms_local;
 940
 941     /* need to clear GPU f output if realloc happened */
 942     if (realloced)
 943     {
 944         nbnxn_ocl_clear_f(nb, nalloc);
 945     }
 946
 947     if (useLjCombRule(nb->nbparam->vdwtype))
 948     {
 949         ocl_copy_H2D_async(d_atdat->lj_comb, nbat->lj_comb, 0,
 950                            natoms*sizeof(cl_float2), ls, bDoTime ? timers->atdat.fetchNextEvent() : nullptr);
 951     }
 952     else
 953     {
 954         ocl_copy_H2D_async(d_atdat->atom_types, nbat->type, 0,
 955                            natoms*sizeof(int), ls, bDoTime ? timers->atdat.fetchNextEvent() : nullptr);
 956
 957     }
 958
 959     if (bDoTime)
 960     {
 961         timers->atdat.closeTimingRegion(ls);
 962     }
 963
 964     /* kick off the tasks enqueued above to ensure concurrency with the search */
 965     cl_error = clFlush(ls);
 966     assert(CL_SUCCESS == cl_error);
 967 }
 968
 969 /*! \brief Releases an OpenCL kernel pointer */
 970 static void free_kernel(cl_kernel *kernel_ptr)
 971 {
 972     cl_int gmx_unused cl_error;
 973
 974     assert(nullptr != kernel_ptr);
 975
 976     if (*kernel_ptr)
 977     {
 978         cl_error = clReleaseKernel(*kernel_ptr);
 979         assert(cl_error == CL_SUCCESS);
 980
 981         *kernel_ptr = nullptr;
 982     }
 983 }
 984
 985 /*! \brief Releases a list of OpenCL kernel pointers */
 986 static void free_kernels(cl_kernel *kernels, int count)
 987 {
 988     int i;
 989
 990     for (i = 0; i < count; i++)
 991     {
 992         free_kernel(kernels + i);
 993     }
 994 }
 995
 996 /*! \brief Free the OpenCL runtime data (context and program).
 997  *
 998  *  The function releases the OpenCL context and program assuciated with the
 999  *  device that the calling PP rank is running on.
1000  *
1001  *  \param runData [in]  porinter to the structure with runtime data.
1002  */
1003 static void free_gpu_device_runtime_data(gmx_device_runtime_data_t *runData)
1004 {
1005     if (runData == nullptr)
1006     {
1007         return;
1008     }
1009
1010     cl_int gmx_unused cl_error;
1011
1012     if (runData->context)
1013     {
1014         cl_error         = clReleaseContext(runData->context);
1015         runData->context = nullptr;
1016         assert(CL_SUCCESS == cl_error);
1017     }
1018
1019     if (runData->program)
1020     {
1021         cl_error         = clReleaseProgram(runData->program);
1022         runData->program = nullptr;
1023         assert(CL_SUCCESS == cl_error);
1024     }
1025
1026 }
1027
1028 //! This function is documented in the header file
1029 void nbnxn_gpu_free(gmx_nbnxn_ocl_t *nb)
1030 {
1031     if (nb == nullptr)
1032     {
1033         return;
1034     }
1035
1036     /* Free kernels */
1037     int kernel_count = sizeof(nb->kernel_ener_noprune_ptr) / sizeof(nb->kernel_ener_noprune_ptr[0][0]);
1038     free_kernels((cl_kernel*)nb->kernel_ener_noprune_ptr, kernel_count);
1039
1040     kernel_count = sizeof(nb->kernel_ener_prune_ptr) / sizeof(nb->kernel_ener_prune_ptr[0][0]);
1041     free_kernels((cl_kernel*)nb->kernel_ener_prune_ptr, kernel_count);
1042
1043     kernel_count = sizeof(nb->kernel_noener_noprune_ptr) / sizeof(nb->kernel_noener_noprune_ptr[0][0]);
1044     free_kernels((cl_kernel*)nb->kernel_noener_noprune_ptr, kernel_count);
1045
1046     kernel_count = sizeof(nb->kernel_noener_prune_ptr) / sizeof(nb->kernel_noener_prune_ptr[0][0]);
1047     free_kernels((cl_kernel*)nb->kernel_noener_prune_ptr, kernel_count);
1048
1049     free_kernel(&(nb->kernel_memset_f));
1050     free_kernel(&(nb->kernel_memset_f2));
1051     free_kernel(&(nb->kernel_memset_f3));
1052     free_kernel(&(nb->kernel_zero_e_fshift));
1053
1054     /* Free atdat */
1055     freeDeviceBuffer(&(nb->atdat->xq));
1056     freeDeviceBuffer(&(nb->atdat->f));
1057     freeDeviceBuffer(&(nb->atdat->e_lj));
1058     freeDeviceBuffer(&(nb->atdat->e_el));
1059     freeDeviceBuffer(&(nb->atdat->fshift));
1060     freeDeviceBuffer(&(nb->atdat->lj_comb));
1061     freeDeviceBuffer(&(nb->atdat->atom_types));
1062     freeDeviceBuffer(&(nb->atdat->shift_vec));
1063     sfree(nb->atdat);
1064
1065     /* Free nbparam */
1066     freeDeviceBuffer(&(nb->nbparam->nbfp_climg2d));
1067     freeDeviceBuffer(&(nb->nbparam->nbfp_comb_climg2d));
1068     freeDeviceBuffer(&(nb->nbparam->coulomb_tab_climg2d));
1069     sfree(nb->nbparam);
1070
1071     /* Free plist */
1072     auto *plist = nb->plist[eintLocal];
1073     freeDeviceBuffer(&plist->sci);
1074     freeDeviceBuffer(&plist->cj4);
1075     freeDeviceBuffer(&plist->imask);
1076     freeDeviceBuffer(&plist->excl);
1077     sfree(plist);
1078     if (nb->bUseTwoStreams)
1079     {
1080         auto *plist_nl = nb->plist[eintNonlocal];
1081         freeDeviceBuffer(&plist_nl->sci);
1082         freeDeviceBuffer(&plist_nl->cj4);
1083         freeDeviceBuffer(&plist_nl->imask);
1084         freeDeviceBuffer(&plist_nl->excl);
1085         sfree(plist_nl);
1086     }
1087
1088     /* Free nbst */
1089     pfree(nb->nbst.e_lj);
1090     nb->nbst.e_lj = nullptr;
1091
1092     pfree(nb->nbst.e_el);
1093     nb->nbst.e_el = nullptr;
1094
1095     pfree(nb->nbst.fshift);
1096     nb->nbst.fshift = nullptr;
1097
1098     /* Free command queues */
1099     clReleaseCommandQueue(nb->stream[eintLocal]);
1100     nb->stream[eintLocal] = nullptr;
1101     if (nb->bUseTwoStreams)
1102     {
1103         clReleaseCommandQueue(nb->stream[eintNonlocal]);
1104         nb->stream[eintNonlocal] = nullptr;
1105     }
1106     /* Free other events */
1107     if (nb->nonlocal_done)
1108     {
1109         clReleaseEvent(nb->nonlocal_done);
1110         nb->nonlocal_done = nullptr;
1111     }
1112     if (nb->misc_ops_and_local_H2D_done)
1113     {
1114         clReleaseEvent(nb->misc_ops_and_local_H2D_done);
1115         nb->misc_ops_and_local_H2D_done = nullptr;
1116     }
1117
1118     free_gpu_device_runtime_data(nb->dev_rundata);
1119     sfree(nb->dev_rundata);
1120
1121     /* Free timers and timings */
1122     delete nb->timers;
1123     sfree(nb->timings);
1124     sfree(nb);
1125
1126     if (debug)
1127     {
1128         fprintf(debug, "Cleaned up OpenCL data structures.\n");
1129     }
1130 }
1131
1132 //! This function is documented in the header file
1133 gmx_wallclock_gpu_nbnxn_t *nbnxn_gpu_get_timings(gmx_nbnxn_ocl_t *nb)
1134 {
1135     return (nb != nullptr && nb->bDoTime) ? nb->timings : nullptr;
1136 }
1137
1138 //! This function is documented in the header file
1139 void nbnxn_gpu_reset_timings(nonbonded_verlet_t* nbv)
1140 {
1141     if (nbv->gpu_nbv && nbv->gpu_nbv->bDoTime)
1142     {
1143         init_timings(nbv->gpu_nbv->timings);
1144     }
1145 }
1146
1147 //! This function is documented in the header file
1148 int nbnxn_gpu_min_ci_balanced(gmx_nbnxn_ocl_t *nb)
1149 {
1150     return nb != nullptr ?
1151            gpu_min_ci_balanced_factor * nb->dev_info->compute_units : 0;
1152 }
1153
1154 //! This function is documented in the header file
1155 gmx_bool nbnxn_gpu_is_kernel_ewald_analytical(const gmx_nbnxn_ocl_t *nb)
1156 {
1157     return ((nb->nbparam->eeltype == eelOclEWALD_ANA) ||
1158             (nb->nbparam->eeltype == eelOclEWALD_ANA_TWIN));
1159 }