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()