Unverified Commit 3c55bded authored by Axel Kohlmeyer's avatar Axel Kohlmeyer Committed by GitHub
Browse files

Merge pull request #1848 from Vsevak/tip4p_gpu_fix 

Fix some issues with lj/cut/tip4p/long/gpu
parents 868df1f6 43c18498

doc/src/pair_lj.rst

+3 −3
Original line number Diff line number Diff line
@@ -90,15 +90,15 @@ pair\_style lj/cut/coul/wolf/omp command 
pair\_style lj/cut/tip4p/cut command

====================================



pair\_style lj/cut/tip4p/cut/gpu command

========================================



pair\_style lj/cut/tip4p/cut/omp command

========================================



pair\_style lj/cut/tip4p/long command

=====================================



pair\_style lj/cut/tip4p/long/gpu command

========================================



pair\_style lj/cut/tip4p/long/omp command

=========================================

lib/gpu/lal_lj_tip4p_long.cpp

+22 −19
Original line number Diff line number Diff line
@@ -198,8 +198,7 @@ void LJTIP4PLongT::loop(const bool _eflag, const bool _vflag) { 
          &nall, &ainum,

          &nbor_pitch, &this->_threads_per_atom,

          &hneight, &m, &TypeO, &TypeH, &alpha,

          &this->atom->q, &tag, &map_array,

          &atom_sametag);

          &this->atom->q);



  GX=static_cast<int>(ceil(static_cast<double>(this->ans->inum())/

                               (BX/this->_threads_per_atom)));
@@ -217,8 +216,7 @@ void LJTIP4PLongT::loop(const bool _eflag, const bool _vflag) { 
          &ainum, &nbor_pitch, &this->_threads_per_atom,

          &hneight, &m, &TypeO, &TypeH, &alpha,

          &this->atom->q, &cutsq, &_qqrd2e, &_g_ewald,

          &cut_coulsq, &cut_coulsqplus, &tag, &map_array,

          &atom_sametag, &this->ansO);

          &cut_coulsq, &cut_coulsqplus, &this->ansO);

  GX=static_cast<int>(ceil(static_cast<double>(this->ans->inum())/BX));

  this->k_pair_distrib.set_size(GX,BX);

  this->k_pair_distrib.run(&this->atom->x, &this->ans->force, &this->ans->engv,
@@ -234,25 +232,30 @@ void LJTIP4PLongT::copy_relations_data(int n, tagint *tag, int *map_array, 
  int nall = n;

  const int hn_sz = n*4; // matrix size = col size * col number

  hneight.resize_ib(hn_sz);

  if (ago == 0)

    hneight.zero();



  m.resize_ib(n);

  m.zero();



  UCL_H_Vec<int> host_tag_write(nall,*(this->ucl_device),UCL_WRITE_ONLY);

  if (ago == 0) {

    hneight.zero();



    {

      UCL_H_Vec<tagint> host_tag_write(nall,*(this->ucl_device),UCL_WRITE_ONLY);

      this->tag.resize_ib(nall);

      for(int i=0; i<nall; ++i) host_tag_write[i] = tag[i];

      ucl_copy(this->tag, host_tag_write, nall, false);

    }



  host_tag_write.resize_ib(max_same);

    UCL_H_Vec<int> host_write(max_same,*(this->ucl_device),UCL_WRITE_ONLY);

    this->atom_sametag.resize_ib(max_same);

  for(int i=0; i<max_same; ++i) host_tag_write[i] = sametag[i];

  ucl_copy(this->atom_sametag, host_tag_write, max_same, false);

    for(int i=0; i<max_same; ++i) host_write[i] = sametag[i];

    ucl_copy(this->atom_sametag, host_write, max_same, false);



  host_tag_write.resize_ib(map_size);

    host_write.resize_ib(map_size);

    this->map_array.resize_ib(map_size);

  for(int i=0; i<map_size; ++i) host_tag_write[i] = map_array[i];

  ucl_copy(this->map_array, host_tag_write, map_size, false);

    for(int i=0; i<map_size; ++i) host_write[i] = map_array[i];

    ucl_copy(this->map_array, host_write, map_size, false);

  }

}

lib/gpu/lal_lj_tip4p_long.cu

+14 −7
Original line number Diff line number Diff line
@@ -16,6 +16,16 @@ 
#ifdef NV_KERNEL



#include "lal_aux_fun1.h"

#ifdef LAMMPS_SMALLBIG

#define tagint int

#endif

#ifdef LAMMPS_BIGBIG

#include "inttypes.h"

#define tagint int64_t

#endif

#ifdef LAMMPS_SMALLSMALL

#define tagint int

#endif

#ifndef _DOUBLE_DOUBLE

texture<float4> pos_tex;

texture<float> q_tex;
@@ -29,7 +39,7 @@ texture<int2> q_tex; 
#define q_tex q_

#endif



ucl_inline int atom_mapping(const __global int *map, int glob) {

ucl_inline int atom_mapping(const __global int *map, tagint glob) {

  return map[glob];

}
@@ -170,7 +180,7 @@ __kernel void k_lj_tip4p_reneigh(const __global numtyp4 *restrict x_, 
    __global int *restrict hneigh,

    __global numtyp4 *restrict m,

    const int typeO, const int typeH,

    const __global int *restrict tag, const __global int *restrict map,

    const __global tagint *restrict tag, const __global int *restrict map,

    const __global int *restrict sametag) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);
@@ -226,9 +236,7 @@ __kernel void k_lj_tip4p_newsite(const __global numtyp4 *restrict x_, 
    __global int *restrict hneigh,

    __global numtyp4 *restrict m,

    const int typeO, const int typeH,

    const numtyp alpha, const __global numtyp *restrict q_,

    const __global int *restrict tag, const __global int *restrict map,

    const __global int *restrict sametag) {

    const numtyp alpha, const __global numtyp *restrict q_) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);

  int i = BLOCK_ID_X*(BLOCK_SIZE_X)+tid;
@@ -268,8 +276,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_, 
    const __global numtyp *restrict cutsq,

    const numtyp qqrd2e, const numtyp g_ewald,

    const numtyp cut_coulsq, const numtyp cut_coulsqplus,

    const __global int *restrict tag, const __global int *restrict map,

    const __global int *restrict sametag, __global acctyp4 *restrict ansO) {

    __global acctyp4 *restrict ansO) {

  int tid, ii, offset;

  atom_info(t_per_atom,ii,tid,offset);

lib/gpu/lal_lj_tip4p_long.h

+1 −1
Original line number Diff line number Diff line
@@ -110,7 +110,7 @@ public: 
  UCL_D_Vec<acctyp4> ansO; // force applied to virtual particle

  // UCL_D_Vec<acctyp4> force_comp;



  UCL_D_Vec<int> tag;

  UCL_D_Vec<tagint> tag;

  UCL_D_Vec<int> map_array;

  UCL_D_Vec<int> atom_sametag;

src/GPU/Install.sh

+2 −2
Original line number Diff line number Diff line
@@ -143,8 +143,8 @@ action pair_ufm_gpu.cpp 
action pair_ufm_gpu.h

action pair_lj_cut_dipole_long_gpu.cpp pair_lj_cut_dipole_long.cpp

action pair_lj_cut_dipole_long_gpu.h pair_lj_cut_dipole_long.cpp

action pair_lj_cut_tip4p_long_gpu.h

action pair_lj_cut_tip4p_long_gpu.cpp

action pair_lj_cut_tip4p_long_gpu.h pair_lj_cut_tip4p_long.cpp

action pair_lj_cut_tip4p_long_gpu.cpp pair_lj_cut_tip4p_long.cpp



# edit 2 Makefile.package files to include/exclude package info

CRA Git | Maintained and supported by SUSTech CRA and CCSE