Commit 43a99362 authored by Vsevak's avatar Vsevak
Browse files

Fix tagint usage

parent d1f3f659
Loading
Loading
Loading
Loading
+14 −10
Original line number Diff line number Diff line
@@ -235,22 +235,26 @@ void LJTIP4PLongT::copy_relations_data(int n, tagint *tag, int *map_array,

  m.resize_ib(n);
  m.zero();

  if (ago == 0) {
    hneight.zero();
    UCL_H_Vec<int> host_tag_write(nall,*(this->ucl_device),UCL_WRITE_ONLY);

    {
      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);
  }
}

+12 −2
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);
+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;