libhoomd/cuda/PotentialExternalGPU.cuh

   1 /*
   2 Highly Optimized Object-oriented Many-particle Dynamics -- Blue Edition
   3 (HOOMD-blue) Open Source Software License Copyright 2009-2014 The Regents of
   4 the University of Michigan All rights reserved.
   5
   6 HOOMD-blue may contain modifications ("Contributions") provided, and to which
   7 copyright is held, by various Contributors who have granted The Regents of the
   8 University of Michigan the right to modify and/or distribute such Contributions.
   9
  10 You may redistribute, use, and create derivate works of HOOMD-blue, in source
  11 and binary forms, provided you abide by the following conditions:
  12
  13 * Redistributions of source code must retain the above copyright notice, this
  14 list of conditions, and the following disclaimer both in the code and
  15 prominently in any materials provided with the distribution.
  16
  17 * Redistributions in binary form must reproduce the above copyright notice, this
  18 list of conditions, and the following disclaimer in the documentation and/or
  19 other materials provided with the distribution.
  20
  21 * All publications and presentations based on HOOMD-blue, including any reports
  22 or published results obtained, in whole or in part, with HOOMD-blue, will
  23 acknowledge its use according to the terms posted at the time of submission on:
  24 http://codeblue.umich.edu/hoomd-blue/citations.html
  25
  26 * Any electronic documents citing HOOMD-Blue will link to the HOOMD-Blue website:
  27 http://codeblue.umich.edu/hoomd-blue/
  28
  29 * Apart from the above required attributions, neither the name of the copyright
  30 holder nor the names of HOOMD-blue's contributors may be used to endorse or
  31 promote products derived from this software without specific prior written
  32 permission.
  33
  34 Disclaimer
  35
  36 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDER AND CONTRIBUTORS ``AS IS'' AND
  37 ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
  38 WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE, AND/OR ANY
  39 WARRANTIES THAT THIS SOFTWARE IS FREE OF INFRINGEMENT ARE DISCLAIMED.
  40
  41 IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
  42 INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
  43 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
  44 DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
  45 LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
  46 OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF
  47 ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  48 */
  49
  50 // Maintainer: jglaser
  51
  52 #include "HOOMDMath.h"
  53 #include "ParticleData.cuh"
  54
  55 #ifdef WIN32
  56 #include <cassert>
  57 #else
  58 #include <assert.h>
  59 #endif
  60
  61 /*! \file PotentialExternalGPU.cuh
  62     \brief Defines templated GPU kernel code for calculating the external forces.
  63 */
  64
  65 #ifndef __POTENTIAL_EXTERNAL_GPU_CUH__
  66 #define __POTENTIAL_EXTERNAL_GPU_CUH__
  67
  68 //! Wraps arguments to gpu_cpef
  69 struct external_potential_args_t
  70     {
  71     //! Construct a external_potential_args_t
  72     external_potential_args_t(Scalar4 *_d_force,
  73               Scalar *_d_virial,
  74               const unsigned int _virial_pitch,
  75               const unsigned int _N,
  76               const Scalar4 *_d_pos,
  77               const BoxDim& _box,
  78               const unsigned int _block_size)
  79                 : d_force(_d_force),
  80                   d_virial(_d_virial),
  81                   virial_pitch(_virial_pitch),
  82                   box(_box),
  83                   N(_N),
  84                   d_pos(_d_pos),
  85                   block_size(_block_size)
  86         {
  87         };
  88
  89     Scalar4 *d_force;                //!< Force to write out
  90     Scalar *d_virial;                //!< Virial to write out
  91     const unsigned int virial_pitch; //!< The pitch of the 2D array of virial matrix elements
  92     const BoxDim& box;         //!< Simulation box in GPU format
  93     const unsigned int N;           //!< Number of particles
  94     const Scalar4 *d_pos;           //!< Device array of particle positions
  95     const unsigned int block_size;  //!< Block size to execute
  96     };
  97
  98 #ifdef NVCC
  99 //! Kernel for calculating external forces
 100 /*! This kernel is called to calculate the external forces on all N particles. Actual evaluation of the potentials and
 101     forces for each particle is handled via the template class \a evaluator.
 102
 103     \param d_force Device memory to write computed forces
 104     \param d_virial Device memory to write computed virials
 105     \param virial_pitch pitch of 2D virial array
 106     \param N number of particles
 107     \param d_pos device array of particle positions
 108     \param box Box dimensions used to implement periodic boundary conditions
 109     \param params per-type array of parameters for the potential
 110
 111 */
 112 template< class evaluator >
 113 __global__ void gpu_compute_external_forces_kernel(Scalar4 *d_force,
 114                                                Scalar *d_virial,
 115                                                const unsigned int virial_pitch,
 116                                                const unsigned int N,
 117                                                const Scalar4 *d_pos,
 118                                                const BoxDim box,
 119                                                const typename evaluator::param_type *params)
 120     {
 121     // start by identifying which particle we are to handle
 122     unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
 123
 124     if (idx >= N)
 125         return;
 126
 127     // read in the position of our particle.
 128     // (MEM TRANSFER: 16 bytes)
 129     Scalar4 posi = d_pos[idx];
 130
 131     // initialize the force to 0
 132     Scalar3 force = make_scalar3(Scalar(0.0), Scalar(0.0), Scalar(0.0));
 133     Scalar virial[6];
 134     for (unsigned int k = 0; k < 6; k++)
 135         virial[k] = Scalar(0.0);
 136     Scalar energy = Scalar(0.0);
 137
 138     unsigned int typei = __scalar_as_int(posi.w);
 139     Scalar3 Xi = make_scalar3(posi.x, posi.y, posi.z);
 140     evaluator eval(Xi, box, params[typei]);
 141
 142     eval.evalForceEnergyAndVirial(force, energy, virial);
 143
 144     // now that the force calculation is complete, write out the result)
 145     d_force[idx].x = force.x;
 146     d_force[idx].y = force.y;
 147     d_force[idx].z = force.z;
 148     d_force[idx].w = energy;
 149
 150     for (unsigned int k = 0; k < 6; k++)
 151         d_virial[k*virial_pitch+idx] = virial[k];
 152     }
 153
 154 //! Kernel driver that computes lj forces on the GPU for LJForceComputeGPU
 155 /*! \param external_potential_args Other arugments to pass onto the kernel
 156     \param d_params Parameters for the potential
 157
 158     This is just a driver function for gpu_compute_external_forces(), see it for details.
 159 */
 160 template< class evaluator >
 161 cudaError_t gpu_compute_external_forces(const external_potential_args_t& external_potential_args,
 162                                     const typename evaluator::param_type *d_params)
 163     {
 164     static unsigned int max_block_size = UINT_MAX;
 165     if (max_block_size == UINT_MAX)
 166         {
 167         cudaFuncAttributes attr;
 168         cudaFuncGetAttributes(&attr, gpu_compute_external_forces_kernel<evaluator>);
 169         max_block_size = attr.maxThreadsPerBlock;
 170         }
 171
 172     unsigned int run_block_size = min(external_potential_args.block_size, max_block_size);
 173
 174     // setup the grid to run the kernel
 175     dim3 grid( external_potential_args.N / run_block_size + 1, 1, 1);
 176     dim3 threads(run_block_size, 1, 1);
 177
 178     // bind the position texture
 179     gpu_compute_external_forces_kernel<evaluator>
 180            <<<grid, threads>>>(external_potential_args.d_force, external_potential_args.d_virial, external_potential_args.virial_pitch, external_potential_args.N, external_potential_args.d_pos, external_potential_args.box, d_params);
 181
 182     return cudaSuccess;
 183     }
 184 #endif
 185
 186 #endif // __POTENTIAL_PAIR_GPU_CUH__