libhoomd/cuda/Enforce2DUpdaterGPU.cu

   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: joaander
  51
  52 #include "Enforce2DUpdaterGPU.cuh"
  53
  54 #ifdef WIN32
  55 #include <cassert>
  56 #else
  57 #include <assert.h>
  58 #endif
  59
  60 #include <stdio.h>
  61
  62 /*! \file Enforce2DUpdaterGPU.cu
  63     \brief Defines GPU kernel code for constraining systems to a 2D plane on
  64     the GPU. Used by Enforce2DUpdaterGPU.
  65 */
  66
  67 //! Constrains partcles to the xy plane on the GPU
  68 /*! \param N number of particles in system
  69     \param d_vel Particle velocities to constrain to xy plane
  70     \param d_accel Particle accelerations to constrain to xy plane
  71 */
  72 extern "C" __global__
  73 void gpu_enforce2d_kernel(const unsigned int N,
  74                           Scalar4 *d_vel,
  75                           Scalar3 *d_accel)
  76     {
  77     int idx = blockIdx.x * blockDim.x + threadIdx.x;
  78
  79     if (idx < N)
  80         {
  81         // read the particle's velocity and acceleration (MEM TRANSFER: 32 bytes)
  82         Scalar4 vel = d_vel[idx];
  83         Scalar3 accel = d_accel[idx];
  84
  85         // zero the z-velocity and z-acceleration(FLOPS: ?)
  86         vel.z = Scalar(0.0);
  87         accel.z = Scalar(0.0);
  88
  89         // write out the results (MEM_TRANSFER: 32 bytes)
  90         d_vel[idx] = vel;
  91         d_accel[idx] = accel;
  92         }
  93     }
  94
  95 /*! \param N number of particles in system
  96     \param d_vel Particle velocities to constrain to xy plane
  97     \param d_accel Particle accelerations to constrain to xy plane
  98 */
  99 cudaError_t gpu_enforce2d(const unsigned int N,
 100                           Scalar4 *d_vel,
 101                           Scalar3 *d_accel)
 102     {
 103     // setup the grid to run the kernel
 104     int block_size = 256;
 105     dim3 grid( (N/block_size) + 1, 1, 1);
 106     dim3 threads(block_size, 1, 1);
 107
 108     // run the kernel
 109     gpu_enforce2d_kernel<<< grid, threads >>>(N, d_vel, d_accel);
 110
 111     return cudaSuccess;
 112     }