From 43a9936241820a1490a9e1f16e20ce1decc4091c Mon Sep 17 00:00:00 2001
From: Vsevak <thevsevak@gmail.com>
Date: Sun, 19 Jan 2020 00:57:45 +0300
Subject: [PATCH] Fix tagint usage

---
 lib/gpu/lal_lj_tip4p_long.cpp | 24 ++++++++++++++----------
 lib/gpu/lal_lj_tip4p_long.cu  | 14 ++++++++++++--
 lib/gpu/lal_lj_tip4p_long.h   |  2 +-
 3 files changed, 27 insertions(+), 13 deletions(-)

diff --git a/lib/gpu/lal_lj_tip4p_long.cpp b/lib/gpu/lal_lj_tip4p_long.cpp
index db94bfe6c3..d44edc8cbd 100644
--- a/lib/gpu/lal_lj_tip4p_long.cpp
+++ b/lib/gpu/lal_lj_tip4p_long.cpp
@@ -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);
-    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<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);
+    }
+
+    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);
   }
 }
 
diff --git a/lib/gpu/lal_lj_tip4p_long.cu b/lib/gpu/lal_lj_tip4p_long.cu
index 9f2f1f6f27..147c460795 100644
--- a/lib/gpu/lal_lj_tip4p_long.cu
+++ b/lib/gpu/lal_lj_tip4p_long.cu
@@ -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);
diff --git a/lib/gpu/lal_lj_tip4p_long.h b/lib/gpu/lal_lj_tip4p_long.h
index 584a53e22b..ab32d1e9f3 100644
--- a/lib/gpu/lal_lj_tip4p_long.h
+++ b/lib/gpu/lal_lj_tip4p_long.h
@@ -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;