← Back to branch summary

~ubuntu-branches/debian/sid/lammps/sid

« back to all changes in this revision

Viewing changes to src/KOKKOS/neigh_full_kokkos.h

  • Committer: Package Import Robot
  • Author(s): Anton Gladky
  • Date: 2015-04-29 23:44:49 UTC
  • mfrom: (5.1.3 experimental)
  • Revision ID: package-import@ubuntu.com-20150429234449-mbhy9utku6hp6oq8
Tags: 0~20150313.gitfa668e1-1
Upload into unstable.

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/KOKKOS/neigh_full_kokkos.h
 
1
/* -*- c++ -*- ----------------------------------------------------------
 
2
   LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
 
3
   http://lammps.sandia.gov, Sandia National Laboratories
 
4
   Steve Plimpton, sjplimp@sandia.gov
 
5
 
 
6
   Copyright (2003) Sandia Corporation.  Under the terms of Contract
 
7
   DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
 
8
   certain rights in this software.  This software is distributed under
 
9
   the GNU General Public License.
 
10
 
 
11
   See the README file in the top-level LAMMPS directory.
 
12
------------------------------------------------------------------------- */
 
13
 
 
14
#include "atom_kokkos.h"
 
15
#include "atom_masks.h"
 
16
#include "domain_kokkos.h"
 
17
 
 
18
using namespace LAMMPS_NS;
 
19
 
 
20
/* ---------------------------------------------------------------------- */
 
21
 
 
22
template<class DeviceType, int HALF_NEIGH>
 
23
void NeighborKokkos::full_bin_kokkos(NeighListKokkos<DeviceType> *list)
 
24
{
 
25
  const int nall = includegroup?atom->nfirst:atom->nlocal;
 
26
  list->grow(nall);
 
27
 
 
28
  NeighborKokkosExecute<DeviceType>
 
29
    data(*list,
 
30
         k_cutneighsq.view<DeviceType>(),
 
31
         k_bincount.view<DeviceType>(),
 
32
         k_bins.view<DeviceType>(),nall,
 
33
         atomKK->k_x.view<DeviceType>(),
 
34
         atomKK->k_type.view<DeviceType>(),
 
35
         atomKK->k_mask.view<DeviceType>(),
 
36
         atomKK->k_molecule.view<DeviceType>(),
 
37
         atomKK->k_tag.view<DeviceType>(),
 
38
         atomKK->k_special.view<DeviceType>(),
 
39
         atomKK->k_nspecial.view<DeviceType>(),
 
40
         atomKK->molecular,
 
41
         nbinx,nbiny,nbinz,mbinx,mbiny,mbinz,mbinxlo,mbinylo,mbinzlo,
 
42
         bininvx,bininvy,bininvz,
 
43
         exclude, nex_type,maxex_type,
 
44
         k_ex1_type.view<DeviceType>(),
 
45
         k_ex2_type.view<DeviceType>(),
 
46
         k_ex_type.view<DeviceType>(),
 
47
         nex_group,maxex_group,
 
48
         k_ex1_group.view<DeviceType>(),
 
49
         k_ex2_group.view<DeviceType>(),
 
50
         k_ex1_bit.view<DeviceType>(),
 
51
         k_ex2_bit.view<DeviceType>(),
 
52
         nex_mol, maxex_mol,
 
53
         k_ex_mol_group.view<DeviceType>(),
 
54
         k_ex_mol_bit.view<DeviceType>(),
 
55
         bboxhi,bboxlo,
 
56
         domain->xperiodic,domain->yperiodic,domain->zperiodic,
 
57
         domain->xprd_half,domain->yprd_half,domain->zprd_half);
 
58
 
 
59
  k_cutneighsq.sync<DeviceType>();
 
60
  k_ex1_type.sync<DeviceType>();
 
61
  k_ex2_type.sync<DeviceType>();
 
62
  k_ex_type.sync<DeviceType>();
 
63
  k_ex1_group.sync<DeviceType>();
 
64
  k_ex2_group.sync<DeviceType>();
 
65
  k_ex1_bit.sync<DeviceType>();
 
66
  k_ex2_bit.sync<DeviceType>();
 
67
  k_ex_mol_group.sync<DeviceType>();
 
68
  k_ex_mol_bit.sync<DeviceType>();
 
69
  atomKK->sync(Device,X_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK|TAG_MASK|SPECIAL_MASK);
 
70
  Kokkos::deep_copy(list->d_stencil,list->h_stencil);
 
71
 
 
72
  data.special_flag[0] = special_flag[0];
 
73
  data.special_flag[1] = special_flag[1];
 
74
  data.special_flag[2] = special_flag[2];
 
75
  data.special_flag[3] = special_flag[3];
 
76
 
 
77
  while(data.h_resize() > 0) {
 
78
    data.h_resize() = 0;
 
79
    deep_copy(data.resize, data.h_resize);
 
80
 
 
81
    MemsetZeroFunctor<DeviceType> f_zero;
 
82
    f_zero.ptr = (void*) k_bincount.view<DeviceType>().ptr_on_device();
 
83
    Kokkos::parallel_for(mbins, f_zero);
 
84
    DeviceType::fence();
 
85
 
 
86
    NeighborKokkosBinAtomsFunctor<DeviceType> f(data);
 
87
 
 
88
    Kokkos::parallel_for(atom->nlocal+atom->nghost, f);
 
89
    DeviceType::fence();
 
90
 
 
91
    deep_copy(data.h_resize, data.resize);
 
92
    if(data.h_resize()) {
 
93
 
 
94
      atoms_per_bin += 16;
 
95
      k_bins = DAT::tdual_int_2d("bins", mbins, atoms_per_bin);
 
96
      data.bins = k_bins.view<DeviceType>();
 
97
      data.c_bins = data.bins;
 
98
    }
 
99
  }
 
100
 
 
101
  if(list->d_neighbors.dimension_0()<nall) {
 
102
    list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", nall*1.1, list->maxneighs);
 
103
    list->d_numneigh = typename ArrayTypes<DeviceType>::t_int_1d("numneigh", nall*1.1);
 
104
    data.neigh_list.d_neighbors = list->d_neighbors;
 
105
    data.neigh_list.d_numneigh = list->d_numneigh;
 
106
  }
 
107
  data.h_resize()=1;
 
108
  while(data.h_resize()) {
 
109
    data.h_new_maxneighs() = list->maxneighs;
 
110
  data.h_resize() = 0;
 
111
 
 
112
  Kokkos::deep_copy(data.resize, data.h_resize);
 
113
  Kokkos::deep_copy(data.new_maxneighs, data.h_new_maxneighs);
 
114
#ifdef KOKKOS_HAVE_CUDA
 
115
    #define BINS_PER_BLOCK 2
 
116
    const int factor = atoms_per_bin<64?2:1;
 
117
    Kokkos::TeamPolicy<DeviceType> config((mbins+factor-1)/factor,atoms_per_bin*factor);
 
118
#else
 
119
    const int factor = 1;
 
120
#endif
 
121
 
 
122
if(newton_pair) {
 
123
  NeighborKokkosBuildFunctor<DeviceType,HALF_NEIGH,1> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
 
124
#ifdef KOKKOS_HAVE_CUDA
 
125
  Kokkos::parallel_for(config, f);
 
126
#else
 
127
  Kokkos::parallel_for(nall, f);
 
128
#endif
 
129
} else {
 
130
  NeighborKokkosBuildFunctor<DeviceType,HALF_NEIGH,0> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
 
131
#ifdef KOKKOS_HAVE_CUDA
 
132
  Kokkos::parallel_for(config, f);
 
133
#else
 
134
  Kokkos::parallel_for(nall, f);
 
135
#endif
 
136
}
 
137
  DeviceType::fence();
 
138
    deep_copy(data.h_resize, data.resize);
 
139
 
 
140
    if(data.h_resize()) {
 
141
      deep_copy(data.h_new_maxneighs, data.new_maxneighs);
 
142
      list->maxneighs = data.h_new_maxneighs() * 1.2;
 
143
      list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs);
 
144
      data.neigh_list.d_neighbors = list->d_neighbors;
 
145
      data.neigh_list.maxneighs = list->maxneighs;
 
146
    }
 
147
  }
 
148
 
 
149
  list->inum = nall;
 
150
  list->gnum = 0;
 
151
 
 
152
}
 
153
 
 
154
/* ---------------------------------------------------------------------- */
 
155
 
 
156
template<class Device>
 
157
KOKKOS_INLINE_FUNCTION
 
158
void NeighborKokkosExecute<Device>::binatomsItem(const int &i) const
 
159
{
 
160
  const int ibin = coord2bin(x(i, 0), x(i, 1), x(i, 2));
 
161
 
 
162
  const int ac = Kokkos::atomic_fetch_add(&bincount[ibin], (int)1);
 
163
  if(ac < bins.dimension_1()) {
 
164
    bins(ibin, ac) = i;
 
165
  } else {
 
166
    resize() = 1;
 
167
  }
 
168
}
 
169
 
 
170
/* ---------------------------------------------------------------------- */
 
171
template<class Device>
 
172
KOKKOS_INLINE_FUNCTION
 
173
int NeighborKokkosExecute<Device>::find_special(const int &i, const int &j) const
 
174
{
 
175
  const int n1 = nspecial(i,0);
 
176
  const int n2 = nspecial(i,1);
 
177
  const int n3 = nspecial(i,2);
 
178
 
 
179
  for (int k = 0; k < n3; k++) {
 
180
    if (special(i,k) == tag(j)) {
 
181
      if (k < n1) {
 
182
        if (special_flag[1] == 0) return -1;
 
183
        else if (special_flag[1] == 1) return 0;
 
184
        else return 1;
 
185
      } else if (k < n2) {
 
186
        if (special_flag[2] == 0) return -1;
 
187
        else if (special_flag[2] == 1) return 0;
 
188
        else return 2;
 
189
      } else {
 
190
        if (special_flag[3] == 0) return -1;
 
191
        else if (special_flag[3] == 1) return 0;
 
192
        else return 3;
 
193
      }
 
194
    }
 
195
  }
 
196
  return 0;
 
197
};
 
198
 
 
199
/* ---------------------------------------------------------------------- */
 
200
 
 
201
template<class Device>
 
202
KOKKOS_INLINE_FUNCTION
 
203
int NeighborKokkosExecute<Device>::exclusion(const int &i,const int &j,
 
204
                                             const int &itype,const int &jtype) const
 
205
{
 
206
  int m;
 
207
 
 
208
  if (nex_type && ex_type(itype,jtype)) return 1;
 
209
 
 
210
  if (nex_group) {
 
211
    for (m = 0; m < nex_group; m++) {
 
212
      if (mask(i) & ex1_bit(m) && mask(j) & ex2_bit(m)) return 1;
 
213
      if (mask(i) & ex2_bit(m) && mask(j) & ex1_bit(m)) return 1;
 
214
    }
 
215
  }
 
216
 
 
217
  if (nex_mol) {
 
218
    for (m = 0; m < nex_mol; m++)
 
219
      if (mask(i) & ex_mol_bit(m) && mask(j) & ex_mol_bit(m) &&
 
220
          molecule(i) == molecule(j)) return 1;
 
221
  }
 
222
 
 
223
  return 0;
 
224
}
 
225
 
 
226
/* ---------------------------------------------------------------------- */
 
227
 
 
228
template<class Device> template<int HalfNeigh,int GhostNewton>
 
229
void NeighborKokkosExecute<Device>::
 
230
   build_Item(const int &i) const
 
231
{
 
232
  /* if necessary, goto next page and add pages */
 
233
  int n = 0;
 
234
  int which = 0;
 
235
  int moltemplate;
 
236
  if (molecular == 2) moltemplate = 1;
 
237
  else moltemplate = 0;
 
238
  // get subview of neighbors of i
 
239
 
 
240
  const AtomNeighbors neighbors_i = neigh_list.get_neighbors(i);
 
241
  const X_FLOAT xtmp = x(i, 0);
 
242
  const X_FLOAT ytmp = x(i, 1);
 
243
  const X_FLOAT ztmp = x(i, 2);
 
244
  const int itype = type(i);
 
245
 
 
246
  const int ibin = coord2bin(xtmp, ytmp, ztmp);
 
247
 
 
248
  const int nstencil = neigh_list.nstencil;
 
249
  const typename ArrayTypes<Device>::t_int_1d_const_um stencil
 
250
    = neigh_list.d_stencil;
 
251
 
 
252
  // loop over all bins in neighborhood (includes ibin)
 
253
  if(HalfNeigh)
 
254
  for(int m = 0; m < c_bincount(ibin); m++) {
 
255
    const int j = c_bins(ibin,m);
 
256
    const int jtype = type(j);
 
257
 
 
258
    //for same bin as atom i skip j if i==j and skip atoms "below and to the left" if using HalfNeighborlists
 
259
    if((j == i) || (HalfNeigh && !GhostNewton && (j < i))  ||
 
260
        (HalfNeigh && GhostNewton && ((j < i) || ((j >= nlocal) &&
 
261
                                       ((x(j, 2) < ztmp) || (x(j, 2) == ztmp && x(j, 1) < ytmp) ||
 
262
                                        (x(j, 2) == ztmp && x(j, 1)  == ytmp && x(j, 0) < xtmp)))))
 
263
      ) continue;
 
264
    if(exclude && exclusion(i,j,itype,jtype)) continue;
 
265
 
 
266
    const X_FLOAT delx = xtmp - x(j, 0);
 
267
    const X_FLOAT dely = ytmp - x(j, 1);
 
268
    const X_FLOAT delz = ztmp - x(j, 2);
 
269
    const X_FLOAT rsq = delx * delx + dely * dely + delz * delz;
 
270
    if(rsq <= cutneighsq(itype,jtype)) {
 
271
      if (molecular) {
 
272
        if (!moltemplate)
 
273
          which = find_special(i,j);
 
274
            /* else if (imol >= 0) */
 
275
            /*   which = find_special(onemols[imol]->special[iatom], */
 
276
            /*                        onemols[imol]->nspecial[iatom], */
 
277
            /*                        tag[j]-tagprev); */
 
278
            /* else which = 0; */
 
279
        if (which == 0){
 
280
          if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
281
        }else if (minimum_image_check(delx,dely,delz)){
 
282
          if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
283
        }
 
284
        else if (which > 0) {
 
285
          if(n<neigh_list.maxneighs) neighbors_i(n++) = j ^ (which << SBBITS);
 
286
        }
 
287
      } else {
 
288
        if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
289
      }
 
290
    }
 
291
  }
 
292
 
 
293
  for(int k = 0; k < nstencil; k++) {
 
294
    const int jbin = ibin + stencil[k];
 
295
    // get subview of jbin
 
296
    if(HalfNeigh&&(ibin==jbin)) continue;
 
297
    //const ArrayTypes<Device>::t_int_1d_const_um =Kokkos::subview<t_int_1d_const_um>(bins,jbin,ALL);
 
298
      for(int m = 0; m < c_bincount(jbin); m++) {
 
299
        const int j = c_bins(jbin,m);
 
300
        const int jtype = type(j);
 
301
 
 
302
        if(HalfNeigh && !GhostNewton && (j < i)) continue;
 
303
        if(!HalfNeigh && j==i) continue;
 
304
        if(exclude && exclusion(i,j,itype,jtype)) continue;
 
305
 
 
306
        const X_FLOAT delx = xtmp - x(j, 0);
 
307
        const X_FLOAT dely = ytmp - x(j, 1);
 
308
        const X_FLOAT delz = ztmp - x(j, 2);
 
309
        const X_FLOAT rsq = delx * delx + dely * dely + delz * delz;
 
310
 
 
311
        if(rsq <= cutneighsq(itype,jtype)) {
 
312
          if (molecular) {
 
313
            if (!moltemplate)
 
314
              which = find_special(i,j);
 
315
            /* else if (imol >= 0) */
 
316
            /*   which = find_special(onemols[imol]->special[iatom], */
 
317
            /*                        onemols[imol]->nspecial[iatom], */
 
318
            /*                        tag[j]-tagprev); */
 
319
            /* else which = 0; */
 
320
            if (which == 0){
 
321
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
322
            }else if (minimum_image_check(delx,dely,delz)){
 
323
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
324
            }
 
325
            else if (which > 0) {
 
326
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j ^ (which << SBBITS);
 
327
            }
 
328
          } else {
 
329
            if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
330
          }
 
331
        }
 
332
 
 
333
      }
 
334
  }
 
335
 
 
336
  neigh_list.d_numneigh(i) = n;
 
337
 
 
338
  if(n >= neigh_list.maxneighs) {
 
339
    resize() = 1;
 
340
 
 
341
    if(n >= new_maxneighs()) new_maxneighs() = n;
 
342
  }
 
343
  neigh_list.d_ilist(i) = i;
 
344
}
 
345
 
 
346
#ifdef KOKKOS_HAVE_CUDA
 
347
extern __shared__ X_FLOAT sharedmem[];
 
348
 
 
349
/* ---------------------------------------------------------------------- */
 
350
 
 
351
template<class DeviceType> template<int HalfNeigh,int GhostNewton>
 
352
__device__ inline
 
353
void NeighborKokkosExecute<DeviceType>::build_ItemCuda(typename Kokkos::TeamPolicy<DeviceType>::member_type dev) const
 
354
{
 
355
  /* loop over atoms in i's bin,
 
356
  */
 
357
  const int atoms_per_bin = c_bins.dimension_1();
 
358
  const int BINS_PER_TEAM = dev.team_size()/atoms_per_bin<1?1:dev.team_size()/atoms_per_bin;
 
359
  const int TEAMS_PER_BIN = atoms_per_bin/dev.team_size()<1?1:atoms_per_bin/dev.team_size();
 
360
  const int MY_BIN = dev.team_rank()/atoms_per_bin;
 
361
 
 
362
  const int ibin = dev.league_rank()*BINS_PER_TEAM+MY_BIN;
 
363
 
 
364
  if(ibin >=c_bincount.dimension_0()) return;
 
365
  X_FLOAT* other_x = sharedmem;
 
366
  other_x = other_x + 5*atoms_per_bin*MY_BIN;
 
367
 
 
368
  int* other_id = (int*) &other_x[4 * atoms_per_bin];
 
369
 
 
370
  int bincount_current = c_bincount[ibin];
 
371
 
 
372
  for(int kk = 0; kk < TEAMS_PER_BIN; kk++) {
 
373
    const int MY_II = dev.team_rank()%atoms_per_bin+kk*dev.team_size();
 
374
  const int i = MY_II < bincount_current ? c_bins(ibin, MY_II) : -1;
 
375
  /* if necessary, goto next page and add pages */
 
376
 
 
377
  int n = 0;
 
378
 
 
379
  X_FLOAT xtmp;
 
380
  X_FLOAT ytmp;
 
381
  X_FLOAT ztmp;
 
382
  int itype;
 
383
  const AtomNeighbors neighbors_i = neigh_list.get_neighbors((i>=0&&i<nlocal)?i:0);
 
384
 
 
385
  if(i >= 0) {
 
386
    xtmp = x(i, 0);
 
387
    ytmp = x(i, 1);
 
388
    ztmp = x(i, 2);
 
389
    itype = type(i);
 
390
    other_x[MY_II] = xtmp;
 
391
    other_x[MY_II + atoms_per_bin] = ytmp;
 
392
    other_x[MY_II + 2 * atoms_per_bin] = ztmp;
 
393
    other_x[MY_II + 3 * atoms_per_bin] = itype;
 
394
  }
 
395
  other_id[MY_II] = i;
 
396
  int test = (__syncthreads_count(i >= 0 && i <= nlocal) == 0);
 
397
 
 
398
  if(test) return;
 
399
 
 
400
  if(i >= 0 && i < nlocal) {
 
401
    #pragma unroll 4
 
402
    for(int m = 0; m < bincount_current; m++) {
 
403
      int j = other_id[m];
 
404
      const int jtype = other_x[m + 3 * atoms_per_bin];
 
405
 
 
406
      //for same bin as atom i skip j if i==j and skip atoms "below and to the left" if using halfneighborlists
 
407
      if((j == i) ||
 
408
         (HalfNeigh && !GhostNewton && (j < i))  ||
 
409
         (HalfNeigh && GhostNewton &&
 
410
            ((j < i) ||
 
411
            ((j >= nlocal) && ((x(j, 2) < ztmp) || (x(j, 2) == ztmp && x(j, 1) < ytmp) ||
 
412
              (x(j, 2) == ztmp && x(j, 1)  == ytmp && x(j, 0) < xtmp)))))
 
413
        ) continue;
 
414
      if(exclude && exclusion(i,j,itype,jtype)) continue;
 
415
      const X_FLOAT delx = xtmp - other_x[m];
 
416
      const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];
 
417
      const X_FLOAT delz = ztmp - other_x[m + 2 * atoms_per_bin];
 
418
      const X_FLOAT rsq = delx * delx + dely * dely + delz * delz;
 
419
 
 
420
      if(rsq <= cutneighsq(itype,jtype)) {
 
421
        if (molecular) {
 
422
          int which = 0;
 
423
          if (!moltemplate)
 
424
            which = find_special(i,j);
 
425
          /* else if (imol >= 0) */
 
426
          /*   which = find_special(onemols[imol]->special[iatom], */
 
427
          /*                        onemols[imol]->nspecial[iatom], */
 
428
          /*                        tag[j]-tagprev); */
 
429
          /* else which = 0; */
 
430
          if (which == 0){
 
431
            if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
432
          }else if (minimum_image_check(delx,dely,delz)){
 
433
            if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
434
          }
 
435
          else if (which > 0) {
 
436
            if(n<neigh_list.maxneighs) neighbors_i(n++) = j ^ (which << SBBITS);
 
437
          }
 
438
        } else {
 
439
          if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
440
        }
 
441
      }
 
442
 
 
443
    }
 
444
  }
 
445
  __syncthreads();
 
446
 
 
447
  const int nstencil = neigh_list.nstencil;
 
448
  const typename ArrayTypes<DeviceType>::t_int_1d_const_um stencil
 
449
    = neigh_list.d_stencil;
 
450
  for(int k = 0; k < nstencil; k++) {
 
451
    const int jbin = ibin + stencil[k];
 
452
 
 
453
    if(ibin == jbin) continue;
 
454
 
 
455
    bincount_current = c_bincount[jbin];
 
456
    int j = MY_II < bincount_current ? c_bins(jbin, MY_II) : -1;
 
457
 
 
458
    if(j >= 0) {
 
459
      other_x[MY_II] = x(j, 0);
 
460
      other_x[MY_II + atoms_per_bin] = x(j, 1);
 
461
      other_x[MY_II + 2 * atoms_per_bin] = x(j, 2);
 
462
      other_x[MY_II + 3 * atoms_per_bin] = type(j);
 
463
     }
 
464
 
 
465
    other_id[MY_II] = j;
 
466
 
 
467
    __syncthreads();
 
468
 
 
469
    if(i >= 0 && i < nlocal) {
 
470
      #pragma unroll 8
 
471
      for(int m = 0; m < bincount_current; m++) {
 
472
        const int j = other_id[m];
 
473
        const int jtype = other_x[m + 3 * atoms_per_bin];
 
474
 
 
475
        //if(HalfNeigh && (j < i))  continue;
 
476
        if(HalfNeigh && !GhostNewton && (j < i)) continue;
 
477
        if(!HalfNeigh && j==i) continue;
 
478
        if(exclude && exclusion(i,j,itype,jtype)) continue;
 
479
 
 
480
        const X_FLOAT delx = xtmp - other_x[m];
 
481
        const X_FLOAT dely = ytmp - other_x[m + atoms_per_bin];
 
482
        const X_FLOAT delz = ztmp - other_x[m + 2 * atoms_per_bin];
 
483
        const X_FLOAT rsq = delx * delx + dely * dely + delz * delz;
 
484
 
 
485
        if(rsq <= cutneighsq(itype,jtype)) {
 
486
          if (molecular) {
 
487
            int which = 0;
 
488
            if (!moltemplate)
 
489
              which = find_special(i,j);
 
490
            /* else if (imol >= 0) */
 
491
            /*   which = find_special(onemols[imol]->special[iatom], */
 
492
            /*                        onemols[imol]->nspecial[iatom], */
 
493
            /*                        tag[j]-tagprev); */
 
494
            /* else which = 0; */
 
495
            if (which == 0){
 
496
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
497
            }else if (minimum_image_check(delx,dely,delz)){
 
498
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
499
            }
 
500
            else if (which > 0) {
 
501
              if(n<neigh_list.maxneighs) neighbors_i(n++) = j ^ (which << SBBITS);
 
502
            }
 
503
          } else {
 
504
            if(n<neigh_list.maxneighs) neighbors_i(n++) = j;
 
505
          }
 
506
        }
 
507
 
 
508
      }
 
509
    }
 
510
    __syncthreads();
 
511
  }
 
512
 
 
513
  if(i >= 0 && i < nlocal) {
 
514
    neigh_list.d_numneigh(i) = n;
 
515
    neigh_list.d_ilist(i) = i;
 
516
  }
 
517
 
 
518
  if(n >= neigh_list.maxneighs) {
 
519
    resize() = 1;
 
520
 
 
521
    if(n >= new_maxneighs()) new_maxneighs() = n;
 
522
  }
 
523
  }
 
524
}
 
525
#endif
 
526
 
 
527
template<class DeviceType>
 
528
void NeighborKokkos::full_bin_cluster_kokkos(NeighListKokkos<DeviceType> *list)
 
529
{
 
530
  const int nall = includegroup?atom->nfirst:atom->nlocal;
 
531
  list->grow(nall);
 
532
 
 
533
  NeighborKokkosExecute<DeviceType>
 
534
    data(*list,
 
535
         k_cutneighsq.view<DeviceType>(),
 
536
         k_bincount.view<DeviceType>(),
 
537
         k_bins.view<DeviceType>(),nall,
 
538
         atomKK->k_x.view<DeviceType>(),
 
539
         atomKK->k_type.view<DeviceType>(),
 
540
         atomKK->k_mask.view<DeviceType>(),
 
541
         atomKK->k_molecule.view<DeviceType>(),
 
542
         atomKK->k_tag.view<DeviceType>(),
 
543
         atomKK->k_special.view<DeviceType>(),
 
544
         atomKK->k_nspecial.view<DeviceType>(),
 
545
         atomKK->molecular,
 
546
         nbinx,nbiny,nbinz,mbinx,mbiny,mbinz,mbinxlo,mbinylo,mbinzlo,
 
547
         bininvx,bininvy,bininvz,
 
548
         exclude, nex_type,maxex_type,
 
549
         k_ex1_type.view<DeviceType>(),
 
550
         k_ex2_type.view<DeviceType>(),
 
551
         k_ex_type.view<DeviceType>(),
 
552
         nex_group,maxex_group,
 
553
         k_ex1_group.view<DeviceType>(),
 
554
         k_ex2_group.view<DeviceType>(),
 
555
         k_ex1_bit.view<DeviceType>(),
 
556
         k_ex2_bit.view<DeviceType>(),
 
557
         nex_mol, maxex_mol,
 
558
         k_ex_mol_group.view<DeviceType>(),
 
559
         k_ex_mol_bit.view<DeviceType>(),
 
560
         bboxhi,bboxlo,
 
561
         domain->xperiodic,domain->yperiodic,domain->zperiodic,
 
562
         domain->xprd_half,domain->yprd_half,domain->zprd_half);
 
563
 
 
564
  k_cutneighsq.sync<DeviceType>();
 
565
  k_ex1_type.sync<DeviceType>();
 
566
  k_ex2_type.sync<DeviceType>();
 
567
  k_ex_type.sync<DeviceType>();
 
568
  k_ex1_group.sync<DeviceType>();
 
569
  k_ex2_group.sync<DeviceType>();
 
570
  k_ex1_bit.sync<DeviceType>();
 
571
  k_ex2_bit.sync<DeviceType>();
 
572
  k_ex_mol_group.sync<DeviceType>();
 
573
  k_ex_mol_bit.sync<DeviceType>();
 
574
 
 
575
  data.special_flag[0] = special_flag[0];
 
576
  data.special_flag[1] = special_flag[1];
 
577
  data.special_flag[2] = special_flag[2];
 
578
  data.special_flag[3] = special_flag[3];
 
579
 
 
580
  atomKK->sync(Device,X_MASK|TYPE_MASK|MASK_MASK|MOLECULE_MASK|TAG_MASK|SPECIAL_MASK);
 
581
  Kokkos::deep_copy(list->d_stencil,list->h_stencil);
 
582
  DeviceType::fence();
 
583
 
 
584
  while(data.h_resize() > 0) {
 
585
    data.h_resize() = 0;
 
586
    deep_copy(data.resize, data.h_resize);
 
587
 
 
588
    MemsetZeroFunctor<DeviceType> f_zero;
 
589
    f_zero.ptr = (void*) k_bincount.view<DeviceType>().ptr_on_device();
 
590
    Kokkos::parallel_for(mbins, f_zero);
 
591
    DeviceType::fence();
 
592
 
 
593
    NeighborKokkosBinAtomsFunctor<DeviceType> f(data);
 
594
 
 
595
    Kokkos::parallel_for(atom->nlocal+atom->nghost, f);
 
596
    DeviceType::fence();
 
597
 
 
598
    deep_copy(data.h_resize, data.resize);
 
599
    if(data.h_resize()) {
 
600
 
 
601
      atoms_per_bin += 16;
 
602
      k_bins = DAT::tdual_int_2d("bins", mbins, atoms_per_bin);
 
603
      data.bins = k_bins.view<DeviceType>();
 
604
      data.c_bins = data.bins;
 
605
    }
 
606
  }
 
607
 
 
608
  if(list->d_neighbors.dimension_0()<nall) {
 
609
    list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", nall*1.1, list->maxneighs);
 
610
    list->d_numneigh = typename ArrayTypes<DeviceType>::t_int_1d("numneigh", nall*1.1);
 
611
    data.neigh_list.d_neighbors = list->d_neighbors;
 
612
    data.neigh_list.d_numneigh = list->d_numneigh;
 
613
  }
 
614
  data.h_resize()=1;
 
615
  while(data.h_resize()) {
 
616
    data.h_new_maxneighs() = list->maxneighs;
 
617
  data.h_resize() = 0;
 
618
 
 
619
  Kokkos::deep_copy(data.resize, data.h_resize);
 
620
  Kokkos::deep_copy(data.new_maxneighs, data.h_new_maxneighs);
 
621
#ifdef KOKKOS_HAVE_CUDA
 
622
    #define BINS_PER_BLOCK 2
 
623
    const int factor = atoms_per_bin<64?2:1;
 
624
    Kokkos::TeamPolicy<DeviceType> config((mbins+factor-1)/factor,atoms_per_bin*factor);
 
625
#else
 
626
    const int factor = 1;
 
627
#endif
 
628
 
 
629
if(newton_pair) {
 
630
  NeighborClusterKokkosBuildFunctor<DeviceType,NeighClusterSize> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
 
631
//#ifdef KOKKOS_HAVE_CUDA
 
632
//  Kokkos::parallel_for(config, f);
 
633
//#else
 
634
  Kokkos::parallel_for(nall, f);
 
635
//#endif
 
636
} else {
 
637
  NeighborClusterKokkosBuildFunctor<DeviceType,NeighClusterSize> f(data,atoms_per_bin * 5 * sizeof(X_FLOAT) * factor);
 
638
//#ifdef KOKKOS_HAVE_CUDA
 
639
//  Kokkos::parallel_for(config, f);
 
640
//#else
 
641
  Kokkos::parallel_for(nall, f);
 
642
//#endif
 
643
}
 
644
  DeviceType::fence();
 
645
    deep_copy(data.h_resize, data.resize);
 
646
 
 
647
    if(data.h_resize()) {
 
648
      deep_copy(data.h_new_maxneighs, data.new_maxneighs);
 
649
      list->maxneighs = data.h_new_maxneighs() * 1.2;
 
650
      list->d_neighbors = typename ArrayTypes<DeviceType>::t_neighbors_2d("neighbors", list->d_neighbors.dimension_0(), list->maxneighs);
 
651
      data.neigh_list.d_neighbors = list->d_neighbors;
 
652
      data.neigh_list.maxneighs = list->maxneighs;
 
653
    }
 
654
  }
 
655
 
 
656
  list->inum = nall;
 
657
  list->gnum = 0;
 
658
 
 
659
}
 
660
 
 
661
/* ---------------------------------------------------------------------- */
 
662
 
 
663
template<class Device> template<int ClusterSize>
 
664
void NeighborKokkosExecute<Device>::
 
665
   build_cluster_Item(const int &i) const
 
666
{
 
667
  /* if necessary, goto next page and add pages */
 
668
  int n = 0;
 
669
 
 
670
  // get subview of neighbors of i
 
671
 
 
672
  const AtomNeighbors neighbors_i = neigh_list.get_neighbors(i);
 
673
  const X_FLOAT xtmp = x(i, 0);
 
674
  const X_FLOAT ytmp = x(i, 1);
 
675
  const X_FLOAT ztmp = x(i, 2);
 
676
  const int itype = type(i);
 
677
 
 
678
  const int ibin = coord2bin(xtmp, ytmp, ztmp);
 
679
 
 
680
  const int nstencil = neigh_list.nstencil;
 
681
  const typename ArrayTypes<Device>::t_int_1d_const_um stencil
 
682
    = neigh_list.d_stencil;
 
683
 
 
684
  for(int k = 0; k < nstencil; k++) {
 
685
    const int jbin = ibin + stencil[k];
 
686
      for(int m = 0; m < c_bincount(jbin); m++) {
 
687
        const int j = c_bins(jbin,m);
 
688
        bool skip = i == j;
 
689
        for(int k = 0; k< (n<neigh_list.maxneighs?n:neigh_list.maxneighs); k++)
 
690
          if((j-(j%ClusterSize)) == neighbors_i(k)) {skip=true;};//{m += ClusterSize - j&(ClusterSize-1)-1; skip=true;}
 
691
 
 
692
        if(!skip) {
 
693
          const int jtype = type(j);
 
694
 
 
695
          const X_FLOAT delx = xtmp - x(j, 0);
 
696
          const X_FLOAT dely = ytmp - x(j, 1);
 
697
          const X_FLOAT delz = ztmp - x(j, 2);
 
698
          const X_FLOAT rsq = delx * delx + dely * dely + delz * delz;
 
699
 
 
700
          if(rsq <= cutneighsq(itype,jtype)) {
 
701
            if(n<neigh_list.maxneighs) neighbors_i(n) = (j-(j%ClusterSize));
 
702
            n++;
 
703
            //m += ClusterSize - j&(ClusterSize-1)-1;
 
704
          }
 
705
        }
 
706
 
 
707
      }
 
708
  }
 
709
 
 
710
  neigh_list.d_numneigh(i) = n;
 
711
 
 
712
  if(n >= neigh_list.maxneighs) {
 
713
    resize() = 1;
 
714
 
 
715
    if(n >= new_maxneighs()) new_maxneighs() = n;
 
716
  }
 
717
  neigh_list.d_ilist(i) = i;
 
718
}