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