src/gromacs/mdlib/nbnxn_ocl/nbnxn_ocl_kernels.cl

   1 /*
   2  * This file is part of the GROMACS molecular simulation package.
   3  *
   4  * Copyright (c) 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 /* Auxiliary kernels */
  36 __kernel void
  37 memset_f3(__global float3 *buf, const float value, const unsigned int Nbuf)
  38 {
  39     unsigned int tidx = get_global_id(0);
  40     if (tidx < Nbuf)
  41     {
  42         buf[tidx] = value;
  43     }
  44 }
  45
  46 __kernel void
  47 memset_f2(__global float2 *buf, const float value, const unsigned int Nbuf)
  48 {
  49     unsigned int tidx = get_global_id(0);
  50     if (tidx < Nbuf)
  51     {
  52         buf[tidx] = value;
  53     }
  54 }
  55
  56 __kernel void
  57 memset_f(__global float *buf, const float value, const unsigned int Nbuf)
  58 {
  59     unsigned int tidx = get_global_id(0);
  60     if (tidx < Nbuf)
  61     {
  62         buf[tidx] = value;
  63     }
  64 }
  65
  66 /* Very few data */
  67 __kernel void
  68 zero_e_fshift(__global float *fshift, __global float *e_lj, __global float *e_el, const unsigned int Nbuf)
  69 {
  70     unsigned int tidx = get_global_id(0);
  71     if (tidx < Nbuf)
  72     {
  73         fshift[tidx] = 0.0f;
  74     }
  75     if (tidx == 0)
  76     {
  77         *e_lj     = 0.0f;
  78         *e_el     = 0.0f;
  79     }
  80 }
  81
  82 /* Generate pruning kernels. */
  83 #define HAVE_FRESH_LIST 1
  84 #include "nbnxn_ocl_kernel_pruneonly.clh"
  85 #undef HAVE_FRESH_LIST
  86 #include "nbnxn_ocl_kernel_pruneonly.clh"
  87
  88 #if defined GMX_OCL_FASTGEN
  89     #define FLAVOR_LEVEL_GENERATOR "nbnxn_ocl_kernels_fastgen.clh"
  90 #elif defined GMX_OCL_FASTGEN_ADD_TWINCUT
  91     #define FLAVOR_LEVEL_GENERATOR "nbnxn_ocl_kernels_fastgen_add_twincut.clh"
  92 #else
  93     #define FLAVOR_LEVEL_GENERATOR "nbnxn_ocl_kernels.clh"
  94 #endif
  95
  96 /* Top-level kernel generation: will generate through multiple inclusion the
  97  * following flavors for all kernels:
  98  * - force-only output;
  99  * - force and energy output;
 100  * - force-only with pair list pruning;
 101  * - force and energy output with pair list pruning.
 102  */
 103
 104 /** Force only **/
 105 #include FLAVOR_LEVEL_GENERATOR
 106 /** Force & energy **/
 107 #define CALC_ENERGIES
 108 #include FLAVOR_LEVEL_GENERATOR
 109 #undef CALC_ENERGIES
 110
 111 /*** Pair-list pruning kernels ***/
 112 /** Force only **/
 113 #define PRUNE_NBL
 114 #include FLAVOR_LEVEL_GENERATOR
 115 /** Force & energy **/
 116 #define CALC_ENERGIES
 117 #include FLAVOR_LEVEL_GENERATOR
 118 #undef CALC_ENERGIES
 119 #undef PRUNE_NBL