# Z-order curve — how to query all nearest neighbors?

TLDR: How to query all nearest neighbors with the least effort? (Iterating over a bigger neighboring area that's continuous and filtering out elements that are not immediate neighbors?)

Suppose the following situation:

• You have a set of 3D points' discrete-grid locations stored in a 1D array (an OpenCL buffer)
• You've assigned each point its Z-index/Morton code
• You sort the array according the Z-indexes
• You generate a lookup-array containing the beginning and end indices of continuous regions of same-Z-index-value particles in the first array

Next suppose that you have an OpenCL kernel which executes for each individual particle. And in that kernel you need to loop over all nearest neighboring particles – not only grid cells – (including the currently processed particle itself) of the currently processed particle (that could be 9 neighbors for 2D case, or 27 for 3D).

All of the neighboring particles (for a processed particle) are stored in the nearest-neighboring grid cells. Note that the number of particles in one cell can be high (e.g. look at the bottom-right cells of the image below).

Question: How can I iterate through all the neighboring particles (ideally in a continuous manner)? If this couldn't be solved via Z-order curve, what other approaches can I use?

Example pseudocode:

__kernel void process_stuff(...) {
foreach ( neighboring_particles ) {  // <- how to achieve this efficiently?
// perform computation
}
}


Image legend:

• Red line = Z-order curve
• The blue dot = the currently processed particle
• Yellow dots = another particles
• Green region = the current particle neighborhood – a set of particles I need to iterate through
• It sounds like you're looking for space partitioning unless I am misunderstanding your question. Jul 28 '16 at 15:59
• @NilsOleTimm yes, but to clarify the context: I'm trying to implement Position Based Fluids. The authors suggest to use CUDA Particles approach for creating a spatial search structure. I've decided to implement it using the "Building the Grid using Sorting" method. I have all the structures created, but now I need to iterate through all particle's neighbors in one OpenCL kernel, yet I don't know how to do this (efficiently). Jul 28 '16 at 18:07
• I realize it's a year late, but the wikipedia article mentions how to off-set a Z-order index in 2D. May 22 '17 at 17:00