def __init__(self, posarg, group=None, h=1., maxneighs=100, index=None, indvec=None, neigharray=None):
"""Initialize fixed radius neighbor search. Neighbors are searched inside a fixed radius
specified by h, which defines the cell size in which the spatial domain is divided during searching.
The Algorithm uses Z-order indexing and parallel counting sort (see Syncgroup.sort() and ZIndex for
details).
Input:
posarg -- the positional argument, must be a 3d vector, e.g. CLFloat4, CLDouble4 or CLReal4.
group -- the Syncgroup posarg belongs to, if not provided an internal group is created
(self._group).
h -- cell size for the neighbor search. radius = h/2.
maxneighs -- maximum neighbor count per item.
index -- Z-order index. --> see ZIndex
indvec -- spatial index --> see ZIndex
neigharray -- array containing neigbors indices of shape( len(posarg), maxneighs).
If not specified available as self.neighbors.
"""
if not group:
group = Syncgroup(posarg)
if not(posarg in group._vars):
raise ValueError('posarg must be member of group.')
self._posarg = posarg
self._group = group
self._maxneighs = CLScalar(maxneighs)
self._indexkernel = ZIndex(posarg=posarg, h=h, index=index, indvec=indvec)
self._group.add(self._indexkernel._indvec)
self._group.setup_sort(index=self._indexkernel._index)
self._neighbors = None
if neigharray:
neigharray.shape = (self._indexkernel._index._value.size, maxneighs)
self._neighbors = neigharray
else:
self._neighbors = CLInt(self._indexkernel._index._value.size, maxneighs)
self._grid_length = CLScalar(len(self._group._grid))#, name='__grid_length')
self._posarg_length = CLScalar(len(self._posarg))#, name='__posarg_length')
src = Neighsearch._find_kern_str
src = src.replace('$posarg_length', self._posarg_length.name)
src = src.replace('$grid_length', self._grid_length.name)
src = src.replace('$posarg', self._posarg.name)
src = src.replace('$neighbors', self._neighbors.name)
src = src.replace('$ind4', self._indexkernel._indvec.name)
src = src.replace('$grid', self._group._grid.name)
src = src.replace('$maxneighs', self._maxneighs.name)
src = src.replace('$h', self._indexkernel._h.name)
varlist = [self._posarg_length, self._grid_length, self._posarg, self._neighbors, self._indexkernel._indvec, self._group._grid, self._maxneighs, self._indexkernel._h]
self._findkern = CLKernel(pre=ZIndex._il3_16_src, varlist=varlist, src=src, name='find_neighbors')
self._reset_kern = CLTemplateKernel(src=Neighsearch._reset_kern_str, name='reset_neighbors')
self._reset_kern.neighbors = self._neighbors
self._reset_kern.maxneighs = self._maxneighs
self._reset_kern.compile()
class Neighsearch(object):
_find_kern_pre = """
int gid = get_global_id(0) + get_global_id(1)*get_global_size(0) + get_global_id(2)*get_global_size(0)*get_global_size(1);
int gsz = get_global_size(0)*get_global_size(1)*get_global_size(2);
int xind = $ind4[gid].s0;
int yind = $ind4[gid].s1;
int zind = $ind4[gid].s2;
real4 this_pos = $posarg[gid];
int id, xc, yc, zc, start;
start = 0;
int counter = 0;
"""
_find_kern_body = """
id = interleave3_16(xind + xc, yind + yc, zind + zc);
if ((id < $grid_length) && (id >= 0)){
start = $grid[id] - 1;
int len_cell = 0;
if (id == 0){
len_cell = start + 1;
}else{
len_cell = (start + 1) - $grid[id - 1];
}
if ((start >= 0) && (start < $posarg_length)){
for (int i = 0; i < len_cell; i++){
if (counter < $maxneighs){
int other_id = start - i;
if (other_id != gid){
real dist = length(this_pos - $posarg[other_id]);
if (dist < $h){
$neighbors[gid*$maxneighs + counter] = other_id;
counter++;
}
}
}
}
}
}
"""
_find_kern_str = ""
_find_kern_str += _find_kern_pre
for i in range(3):
for j in range(3):
for k in range(3):
_find_kern_str += 'xc = ' + str(i - 1) + ';\n'
_find_kern_str += 'yc = ' + str(j - 1) + ';\n'
_find_kern_str += 'zc = ' + str(k - 1) + ';\n\n'
_find_kern_str += _find_kern_body
_reset_kern_str = """
int gid = get_global_id(0) + get_global_id(1)*get_global_size(0) + get_global_id(2)*get_global_size(0)*get_global_size(1);
$neighbors$[gid] = -1;
"""
def __init__(self, posarg, group=None, h=1., maxneighs=100, index=None, indvec=None, neigharray=None):
"""Initialize fixed radius neighbor search. Neighbors are searched inside a fixed radius
specified by h, which defines the cell size in which the spatial domain is divided during searching.
The Algorithm uses Z-order indexing and parallel counting sort (see Syncgroup.sort() and ZIndex for
details).
Input:
posarg -- the positional argument, must be a 3d vector, e.g. CLFloat4, CLDouble4 or CLReal4.
group -- the Syncgroup posarg belongs to, if not provided an internal group is created
(self._group).
h -- cell size for the neighbor search. radius = h/2.
maxneighs -- maximum neighbor count per item.
index -- Z-order index. --> see ZIndex
indvec -- spatial index --> see ZIndex
neigharray -- array containing neigbors indices of shape( len(posarg), maxneighs).
If not specified available as self.neighbors.
"""
if not group:
group = Syncgroup(posarg)
if not(posarg in group._vars):
raise ValueError('posarg must be member of group.')
self._posarg = posarg
self._group = group
self._maxneighs = CLScalar(maxneighs)
self._indexkernel = ZIndex(posarg=posarg, h=h, index=index, indvec=indvec)
self._group.add(self._indexkernel._indvec)
self._group.setup_sort(index=self._indexkernel._index)
self._neighbors = None
if neigharray:
neigharray.shape = (self._indexkernel._index._value.size, maxneighs)
self._neighbors = neigharray
else:
self._neighbors = CLInt(self._indexkernel._index._value.size, maxneighs)
self._grid_length = CLScalar(len(self._group._grid))#, name='__grid_length')
self._posarg_length = CLScalar(len(self._posarg))#, name='__posarg_length')
src = Neighsearch._find_kern_str
src = src.replace('$posarg_length', self._posarg_length.name)
src = src.replace('$grid_length', self._grid_length.name)
src = src.replace('$posarg', self._posarg.name)
src = src.replace('$neighbors', self._neighbors.name)
src = src.replace('$ind4', self._indexkernel._indvec.name)
src = src.replace('$grid', self._group._grid.name)
src = src.replace('$maxneighs', self._maxneighs.name)
src = src.replace('$h', self._indexkernel._h.name)
varlist = [self._posarg_length, self._grid_length, self._posarg, self._neighbors, self._indexkernel._indvec, self._group._grid, self._maxneighs, self._indexkernel._h]
self._findkern = CLKernel(pre=ZIndex._il3_16_src, varlist=varlist, src=src, name='find_neighbors')
self._reset_kern = CLTemplateKernel(src=Neighsearch._reset_kern_str, name='reset_neighbors')
self._reset_kern.neighbors = self._neighbors
self._reset_kern.maxneighs = self._maxneighs
self._reset_kern.compile()
@property
def neighbors(self):
return self._neighbors
@timing
def search(self, h=None, maxneighs=None):
self._indexkernel(h=h)
if maxneighs:
self._maxneighs.value = maxneighs
self._neighbors.set_shape_wo_read( (self._indexkernel._index._value.size, maxneighs) )
elif not(self._indexkernel._index._value.size == self._neighbors.shape[0]):
self._neighbors.set_shape_wo_read( (self._indexkernel._index._value.size, self._maxneighs.value) )
self._reset_kern()
self._group.sort()
self._grid_length.value = len(self._group._grid)
self._posarg_length.value = len(self._posarg)
self._findkern()
def info(self):
return self._indexkernel.info()
