src/gromacs/nbnxm/opencl/nbnxm_ocl_kernels.cl

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 2017,2018,2019, 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 /* Auxiliary kernels */
  36
  37 /* Very few data */
  38 __kernel void
  39 zero_e_fshift(__global float *fshift, __global float *e_lj, __global float *e_el, const unsigned int Nbuf)
  40 {
  41     unsigned int tidx = get_global_id(0);
  42     if (tidx < Nbuf)
  43     {
  44         fshift[tidx] = 0.0f;
  45     }
  46     if (tidx == 0)
  47     {
  48         *e_lj     = 0.0f;
  49         *e_el     = 0.0f;
  50     }
  51 }
  52
  53 /* Generate pruning kernels. */
  54 #define HAVE_FRESH_LIST 1
  55 #include "nbnxm_ocl_kernel_pruneonly.clh"
  56 #undef HAVE_FRESH_LIST
  57 #include "nbnxm_ocl_kernel_pruneonly.clh"
  58
  59 #if defined GMX_OCL_FASTGEN
  60     #define FLAVOR_LEVEL_GENERATOR "nbnxm_ocl_kernels_fastgen.clh"
  61 #elif defined GMX_OCL_FASTGEN_ADD_TWINCUT
  62     #define FLAVOR_LEVEL_GENERATOR "nbnxm_ocl_kernels_fastgen_add_twincut.clh"
  63 #else
  64     #define FLAVOR_LEVEL_GENERATOR "nbnxm_ocl_kernels.clh"
  65 #endif
  66
  67 /* Top-level kernel generation: will generate through multiple inclusion the
  68  * following flavors for all kernels:
  69  * - force-only output;
  70  * - force and energy output;
  71  * - force-only with pair list pruning;
  72  * - force and energy output with pair list pruning.
  73  */
  74
  75 /** Force only **/
  76 #include FLAVOR_LEVEL_GENERATOR
  77 /** Force & energy **/
  78 #define CALC_ENERGIES
  79 #include FLAVOR_LEVEL_GENERATOR
  80 #undef CALC_ENERGIES
  81
  82 /*** Pair-list pruning kernels ***/
  83 /** Force only **/
  84 #define PRUNE_NBL
  85 #include FLAVOR_LEVEL_GENERATOR
  86 /** Force & energy **/
  87 #define CALC_ENERGIES
  88 #include FLAVOR_LEVEL_GENERATOR
  89 #undef CALC_ENERGIES
  90 #undef PRUNE_NBL