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