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