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