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