1 // **************************************************************************
2 //                                  atom.cu
3 //                             -------------------
4 //                           W. Michael Brown (ORNL)
5 //
6 //  Device code for handling CPU generated neighbor lists
7 //
8 // __________________________________________________________________________
9 //    This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)
10 // __________________________________________________________________________
11 //
12 //    begin                :
13 //    email                : brownw@ornl.gov
14 // ***************************************************************************/
15 
16 #ifdef NV_KERNEL
17 #include "lal_preprocessor.h"
18 #endif
19 
kernel_unpack(__global int * dev_nbor,const __global int * dev_ij,const int inum,const int t_per_atom)20 __kernel void kernel_unpack(__global int *dev_nbor,
21                             const __global int *dev_ij,
22                             const int inum, const int t_per_atom) {
23   int tid=THREAD_ID_X;
24   int offset=tid & (t_per_atom-1);
25   int ii=fast_mul((int)BLOCK_ID_X,(int)(BLOCK_SIZE_X)/t_per_atom)+tid/t_per_atom;
26 
27   if (ii<inum) {
28     __global int *nbor=dev_nbor+ii+inum;
29     int numj=*nbor;
30     nbor+=inum;
31     const __global int *list=dev_ij+*nbor;
32     const __global int *list_end=list+numj;
33     list+=offset;
34     nbor+=fast_mul(ii,t_per_atom-1)+offset;
35     int stride=fast_mul(t_per_atom,inum);
36 
37     for ( ; list<list_end; list++) {
38       *nbor=*list;
39       nbor+=stride;
40     }
41   } // if ii
42 }
43 
44