Randomizing OpenCL/CUDA indexing by some permutation

Typically I work on the float volume of the integer sizes vdimx, vdimy, vdimz in OpenCL. On the beginning of the kernel, I typically have the code such as

int i = get_global_id(0);
int j = get_global_id(1);
int k = get_global_id(2);

It seems that OpenCL runtime allocates i,j,k very close to each other at given time. So if vdimx=vdimy=vdimz=1024 it is likely that the first indices run will be (0,0,0), (0,0,1) (0,0,2) and not (123,5,951), (350,3,221), (1,1021,151). But in certain applications, e.g. CT projector, closeness of the voxels causes that they project to the same pixels and in turn it causes atomic operations to be slow as multiple GPU threads are trying to write to the same position in memory.

I would like to avoid this to have e.g.

int i = get_global_id(0);
int j = get_global_id(1);
int k = get_global_id(2);
i = perm(i);
j = perm(j);
k = perm(k);

to randomize positions of the voxels. Now the question is how to implement the permutation efficiently and correctly. E.g. if I do

#define perm(i) ((i * 10) % vdimx);

will I obtain 1-1 mapping for any vdimx? When I multiply by prime number which does not divide vidimx do I obtain permutation? Is there some other technique to randomize indices?

Not exactly an answer to the posed question, but hopefully an answer to the underlying problem: dealing with conflicts.

When you access your memory, all threads in a warp need to wait until data for all threads in the warp are fetched. If all the data is physically next to each other, that is resolved in one or two memory transactions. But if your data is located at random places, all threads need to wait for 32 memory transactions to complete (not to mention cache trashing).

That is why I would advise you against randomizing your indices. Instead, focus your efforts to make writing the results faster. Try to resolve the atomic conflicts on the warp/block level, before committing to global memory.

For example: each thread in a block knows into which cell it wants to write to. It then needs to compare its own cell index with the index that other threads in a warp or block want to write to. If they are the same, a winner should be chosen somehow (e.g. a thread with a lower threadIdx wins). Then, only the winners do an atomic operation on global.

Picking winners among n threads in a block is something that (I believe) can be done in log-n steps with reduction algorithms, using only registers and shared memory and will be faster than any subsequent atomic operation on global. I am not sure about the most efficient approach, but a more straightforward is to simply sort all values by index, and then each thread checks if its predecessor has the same index.

I just found an example of a warp-level bitonic sort over here: https://tschmidt23.github.io/cse599i/CSE%20599%20I%20Accelerated%20Computing%20-%20Programming%20GPUs%20Lecture%2018.pdf

If you know a better approach for choosing the winners, please do share! There must be some articles on the topic, even if not CUDA-specific.