Page Menu
Home
c4science
Search
Configure Global Search
Log In
Files
F79466531
lal_neighbor_cpu.cu
No One
Temporary
Actions
Download File
Edit File
Delete File
View Transforms
Subscribe
Mute Notifications
Award Token
Subscribers
None
File Metadata
Details
File Info
Storage
Attached
Created
Mon, Aug 26, 13:29
Size
1 KB
Mime Type
text/x-c
Expires
Wed, Aug 28, 13:29 (1 d, 23 h)
Engine
blob
Format
Raw Data
Handle
20139525
Attached To
rLAMMPS lammps
lal_neighbor_cpu.cu
View Options
// **************************************************************************
// 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
// ***************************************************************************/
#ifdef NV_KERNEL
#include "lal_preprocessor.h"
#endif
__kernel void kernel_unpack(__global int *dev_nbor, __global int *dev_ij,
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) {
__global int *nbor=dev_nbor+ii+inum;
int numj=*nbor;
nbor+=inum;
__global int *list=dev_ij+*nbor;
__global 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++) {
*nbor=*list;
nbor+=stride;
}
} // if ii
}
Event Timeline
Log In to Comment